Skip to content
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.

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

2 issues here #277

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

2 issues here #277

wants to merge 227 commits into from

Conversation

eric1hello
Copy link

I found 2 issues here:
---- shader.cc
I think the pI1 shall be pI2.
if ((pI1->oprnd_type == INT_OP) || (pI1->oprnd_type == UN_OP)) { //these counters get added up in mcPat to compute scheduler power
m_stats->m_num_INTdecoded_insn[m_sid]++;
---- I tried to enable cooperative_groups in bellow cuda, but seem it doesn't work , something issue with PTX, do you know the reasons?

device int atomicAggInc(int *ptr) {
auto g = cg::coalesced_threads();
int prev;

if (g.thread_rank() == 0)
prev = atomicAdd(ptr,g.size());

prev = g.thread_rank() + g.shfl(prev,0);

return prev;

}

global void vectorAdd(float *A, const float *B, float *C , int numElements) {

int i = blockDim.x * blockIdx.x + threadIdx.x;

//if (i < numElements) {
// C[i] = A[i] + B[i];
//}
if ( i%10 == 0){
int rankIdx = atomicAggInc(&count);
printf ("blockIdx = %d, threadIdx = %d, rank = %d \n",blockIdx.x ,threadIdx.x,rankIdx);
}
}

gvoskuilen and others added 30 commits October 26, 2020 15:19
Fix for bugs in lazy write handling
change address type into ull
do not truncate 32 MSB bits of the memory address
Fix cache hash function and renaming
JRPan and others added 30 commits January 10, 2024 14:13
Adding a non-zero return on error
* migrate_cmake: add package dependency checking

* migrate_cmake: port setup_environment to CMake

* migrate_cmake: break dependency checking and env export gen to different .cmake files

* migrate_cmake: use CUDAToolkit_FOUND to test for CUDA compiler

* migrate_cmake: use CUDAToolkit_FOUND to test for CUDA compiler

* migrate_cmake: use CUDAToolkit_FOUND to test for CUDA compiler

* migrate_cmake: properly parse for cuda version number

* migrate_cmake: set highest CUDA supported to be 11.10.x

* migrate_cmake: specify top level CMake file

* migrate_cmake: add libcuda cmake file

* migrate_cmake: use global compiler options and definitions

* migrate_cmake: add cmake file to src

* migrate_cmake: add cmake files for cuda-sim folder

* migrate_cmake: add cmake files to gpgpu-sim folder

* migrate_cmake: add cmake files for intersim

* migrate_cmake: add short test using cmake

* migrate_cmake: bump CXX standard requirement to 17

* Add cmake files for accelwattch

* migrate_cmake: remove use of GLOB to grab source files

* migrate_cmake: comment out the write protection on generated instructions.h

* migrate_cmake: create sym folder and add newline to generated setup file

* migrate_cmake: fix some path issues

* migrate_cmake: let cmake thinks flex and bison generate CXX files

* migrate_cmake: fix not linking pthread properly

* migrate_cmake: remove debug message

* migrate_cmake: add empty libopencl cmake file

* migrate_cmake: install phase and runtime version detect

	* Added install phase to install the shared object
	  and add symlinks

	* Changes with CUDA toolkit will be detected and
	  triggered a rebuild

	* GPGPU-Sim detailed version string will be updated
	  on each build

* Typo fix and fix correct bin dir

* Replace gcc -> g++ in intersim

* ignore setup

* check CMAKE_BUILD_TYPE

* set DCMAKE_BUILD_TYPE

---------

Co-authored-by: JRPAN <25518778+JRPan@users.noreply.github.com>
CMAKE_BUILD_TYPE should be inside ${}
LDGSTS/LDGDEPBAR was introduced #62, but it's increment part was deleted
by mistake. So add it.

In some applications, ldgsts may not exist between ldgdepbar. In
such cases, add exception handling logic to insert an empty vector.

Reported-by: Okkyun Woo <okkyun.w@postech.ac.kr>

Signed-off-by: Wonhyuk Yang <wonhyuk@postech.ac.kr>
* remove implicit casting, cleanup unused bank_warp_shift parameter

* update cu init function prototype

* remove m_bank_warp_shift from function call
* add automated clang formatter

* Automated clang-format

* use /bin/bash and add print

* use default checkout ref

* Format only after tests are success

* Run CI on merge group

---------

Co-authored-by: barnes88 <barnes88@users.noreply.github.com>
Co-authored-by: JRPAN <25518778+JRPan@users.noreply.github.com>
* run formatter only on PR

* remove unused & unintilized variables

* fix signed & unsigned comparison warning

* enable merge queue

* resolve conflict

* in formatter, checkout the forked repo, not the base repo in PR

* Try to use jenkins for formatter

* Automated Format

---------

Co-authored-by: purdue-jenkins <purdue-jenkins@users.noreply.github.com>
* Temp commit for Justin and Cassie to sync on
code changes for adding per-stream status.

* Resolved compile errors.

* Removed redundant parameter

* Passed cuda_stream_id from accelsim to gpgpusim

* Cleaned up unused changes

* Changed vector to map, having operator problems.

* StreamID defaults to zero

* Implemented streams to inc_stats and so on

* Fixed TOTAL_ACCESS counts

* Implemented GLOBAL_TIMER.

* Fixed m_shader->get_kernel SEGFAULT issue in shader.cc.

* Use warp_init to track streamID instead of issue_warp

* Removed temp debug print

* Modified cache_stats to only print data from latest finished stream

Added optional arg to cache_stats::print_stats, cache_stats::print_fail_stats and their upstream functions. When streamID is specified, print stats
from that stream. When not specified, print all stats.

NOTE: current implementation depending on streamid never equals -1

* Removed default arg values of streamID

* modified constructor of mem_fetch to pass in streamID

* changed get_streamid to get_streamID

* Added TODO to gpgpusim_entrypoint.cc and power_stat.cc

* Only collect power stats when enabled

* print last finished stream in PTX mode using last_streamID

* take out additional printf

* Add a field to baseline cache to indicate cache level

* save gpu object in cache

* Print stream ID only once per kernel

* rm test print

* use -1 for default stream id

* cleanup debug prints

* remove GLOABL_TIMER

* Automated clang-format

* Should be correct to print everything in power model

* addressing concerns & errors

* Automated clang-format

* add m_stats_pw in operator+

* Automated Format

---------

Co-authored-by: Justin Qiao <sqiao6@wisc.edu>
Co-authored-by: Justin Qiao <71228724+ShichenQiao@users.noreply.github.com>
Co-authored-by: Tim Rogers <timrogers@gmail.com>
Co-authored-by: JRPan <JRPan@users.noreply.github.com>
Co-authored-by: purdue-jenkins <purdue-jenkins@users.noreply.github.com>
* we have gcc-11 now. Check version for more than 2 digits.

* version detection as well - And support c++ 11 by default
* Add accommodations to run gpgpusim with SST simulation framework through balar

* Output setup_environment options when sourcing

* Add SST directive check when creating sim thread

* Add sst side test for jenkins

* sst-integration: update Jenkinsfile with offical sst-elements repo and fix bugs in pipeline script

* sst-integration: direct jenkins to rebuild gpgpusim before testing for sst

* sst-integration: fix bugs in sst repos config

* sst-integration: let Jenkins rebuilds simulator

Since the simulator needs to be configured with both normal mode and sst mode, need to rebuild make target to clean prior runs.

* sst-integration: Update Jenkinsfile to source env vars when running balar test

* sst-integration: refactor code to remove __SST__ flag

* sst-integration: fix a bug that init cluster twice for sst

* sst-integration: fix a bug of not sending mem packets to SST

* sst-integration: remove sst flags from makefiles and setup_env

* sst-integration: add comments to SST changes

* sst-integration: remove rebuilding simulator in jenkins when testing for SST

* sst-integration: revert simulator build script

* Add a function to support querying function argument info for SST

* sst-integration: add version detection for vanadis binary

* Automated Format

* add version detection support for gcc 10+

* sst-integration: add cudaMallocHost for SST

* sst-integration: fix a compilation bug

* sst-integration: add sst balar unittest CI

* sst-integration: specify GPU_ARCH for CI test

* sst-integration: use bash for github actions

* sst-integration: use https links for sst repos

* sst-integration: add SST dependencies to CI config

* sst-integration: remove sudo

* sst-integration: default to yes for apt install

* sst-integration: add manual trigger for github action

* sst-integration: remove wrong on event

* sst-integration: limit CPU usage for compilation

* sst-integration: fix wrong path

* sst-integration: use personal repo for testing

* sst-integration: remove sst-core source in CI to free space

* sst-integration: SST_Cycle use print stats with stream id

* Automated Format

* sst-integration: check for diskspace and try to clean it

* sst-integration: move out of docker image

* sst-integration: testing for ci path

* sst-integration: fix syntax

* sst-integration: pass env vars

* sst-integration: set env properly

* sst-integration: merge LLVM build and test into same job

* sst-integration: fix step order

* sst-integration: checkout correct branch for env-setup

* sst-integration: remove resourcing gpu apps

* sst-integration: revert back to docker github action

* sst-integration: enable debug trace for sst testing

* sst-integration: resourcing gpu app for env vars

* sst-integration: use GPUAPPS_ROOT for path for gpu app

* sst-integration: use GPUAPPS_ROOT for path for gpu app

* sst-integration: enable parallel ci tests and fix not returning with cudaMallocHostSST

* sst-integration: using debug flag for CI run

* sst-integration: revert debug ci run

* sst-integration: CI skips cuda sdk download and launch multiple jobs

* sst-integration: reenable parallel tests

* sst-integration: reduce concurrent test thread count

* sst-integration: skip long test for github runner

* sst-integration: try running CI with single core

* sst-integrtion: add callback to SST to check thread sync is done in SST_Cycle()

* sst-integration: ignore lookup if already found and add callbacks to SST

* Automated Format

* sst-integration: add support for indirect texture access

* Automated Format

* sste-integration: fix up for PR

* Automated Format

---------

Co-authored-by: purdue-jenkins <purdue-jenkins@users.noreply.github.com>
* fix_sst_callbacks: add weak definitions for sst callbacks

* Automated Format

---------

Co-authored-by: purdue-jenkins <purdue-jenkins@users.noreply.github.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.