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

Change addresses to use 64 bits. Fixes errors with generic addressing. #200

Open
wants to merge 1 commit into
base: dev
Choose a base branch
from
Open

Conversation

AKKamath
Copy link

@AKKamath AKKamath commented Sep 23, 2020

Fixes errors due to address overflow.
Below is an example of correct code that would throw an error in GPGPU-Sim:

#include <iostream>
using namespace std;
#define VALUE 256ULL * 1024ULL * 1024ULL
__global__ void test(int *data1, int *data2)
{
    int value;
    asm volatile("ld.s32 %0, [%1];": "=r"(value) : "l"(data1));
    asm volatile("st.s32 [%0], %1;":: "l"(data2), "r"(value));
}

int main()
{
    int *data1, *data2;
    cudaMalloc((void**)&data1, VALUE * sizeof(int));
    cudaMalloc((void**)&data2, VALUE * sizeof(int));
    int dummy = 100;
    cudaMemcpy(&data1[VALUE - 1], &dummy, sizeof(int), cudaMemcpyHostToDevice);
    test<<<1, 1>>>(&data1[VALUE - 1], data2);
    int data;
    cudaMemcpy(&data, data2, sizeof(int), cudaMemcpyDeviceToHost);
    cout << data << "\n";
}

The address for global memory starts at 0xc000000. By allocating VALUE * sizeof(int) size and using the last index, this address overflows a 32 bit location, causing the simulator to incorrectly read the address.

@AKKamath
Copy link
Author

Fixes Issue #71

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.

1 participant