Skip to content

[IR] Add unroll pragma for loop statement#180

Merged
yaoyaoding merged 2 commits intohidet-org:mainfrom
yaoyaoding:add-unroll
Apr 18, 2023
Merged

[IR] Add unroll pragma for loop statement#180
yaoyaoding merged 2 commits intohidet-org:mainfrom
yaoyaoding:add-unroll

Conversation

@yaoyaoding
Copy link
Copy Markdown
Member

Example:

    from hidet.lang import printf, attr, grid, repeat

    with hidet.script_module() as script_module:
        @hidet.script
        def example():
            attr.func_kind = 'host_kernel'

            for i in grid(10, unroll=True):
                printf("i = %d\n", i)

            for i in grid(10, unroll=2):
                printf("i = %d\n", i)

            for i, j in grid(2, 5, unroll=[True, None]):
                printf("i = %d, j = %d\n", i, j)

            for w in range(32):
                for i, j in repeat(2, 8, unroll=[True, None]).spatial(4, 8).on(w):
                    printf("i = %d, j = %d\n", i, j)

would generate:

__host__ void hidet_example() {
  #pragma unroll
  for (int32_t i = 0; (i < 10); i = (i + 1)) {
    printf("i = %d\n", i);
  } 
  #pragma unroll 2
  for (int32_t i_1 = 0; (i_1 < 10); i_1 = (i_1 + 1)) {
    printf("i = %d\n", i_1);
  } 
  #pragma unroll
  for (int32_t i_2 = 0; (i_2 < 2); i_2 = (i_2 + 1)) {
    for (int32_t j = 0; (j < 5); j = (j + 1)) {
      printf("i = %d, j = %d\n", i_2, j);
    } 
  } 
  for (int32_t w = 0; (w < 32); w = (w + 1)) {
    #pragma unroll
    for (int32_t i_3 = 0; (i_3 < 2); i_3 = (i_3 + 1)) {
      for (int32_t i_4 = 0; (i_4 < 8); i_4 = (i_4 + 1)) {
        printf("i = %d, j = %d\n", ((i_3 * 4) + (w / 8)), ((i_4 * 8) + (w % 8)));
      } 
    } 
  } 
}

@yaoyaoding yaoyaoding merged commit 48e57cb into hidet-org:main Apr 18, 2023
@yaoyaoding yaoyaoding deleted the add-unroll branch April 18, 2023 06:20
vadiklyutiy added a commit that referenced this pull request Jul 22, 2024
We did not limit the size of `multithreading.Pool` by number on jobs
that cause overhead on process creation.

With this PR take a minimum between CPU count and number of jobs. 

**Results**
`time python bench_vision.py resnet50 --params 1x3x224x224 --dtype
float16 --backend hidet`
| Name| Before | After | Improvement|
|--------|--------|--------|--------|
| resnet50 | 11m56  | 11m13 | 6.4% |
vadiklyutiy added a commit that referenced this pull request Jul 23, 2024
We did not limit the size of `multithreading.Pool` by number on jobs
that cause overhead on process creation.

With this PR take a minimum between CPU count and number of jobs. 

**Results**
`time python bench_vision.py resnet50 --params 1x3x224x224 --dtype
float16 --backend hidet`
| Name| Before | After | Improvement|
|--------|--------|--------|--------|
| resnet50 | 11m56  | 11m13 | 6.4% |
vadiklyutiy added a commit that referenced this pull request Dec 26, 2024
We did not limit the size of `multithreading.Pool` by number on jobs
that cause overhead on process creation.

With this PR take a minimum between CPU count and number of jobs. 

**Results**
`time python bench_vision.py resnet50 --params 1x3x224x224 --dtype
float16 --backend hidet`
| Name| Before | After | Improvement|
|--------|--------|--------|--------|
| resnet50 | 11m56  | 11m13 | 6.4% |
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