Skip to content

[Clang][GPU] Fix unit test for NVPTX tid.x intrinsic #136297

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

Merged
merged 1 commit into from
Apr 18, 2025

Conversation

jurahul
Copy link
Contributor

@jurahul jurahul commented Apr 18, 2025

  • llvm.nvvm.read.ptx.sreg.tid.x does not have the result range attribute yet.

- llvm.nvvm.read.ptx.sreg.tid.x does not have the result range
  attribute yet.
@jurahul jurahul marked this pull request as ready for review April 18, 2025 11:59
@llvmbot llvmbot added the clang Clang issues not falling into any other category label Apr 18, 2025
@llvmbot
Copy link
Member

llvmbot commented Apr 18, 2025

@llvm/pr-subscribers-clang

Author: Rahul Joshi (jurahul)

Changes
  • llvm.nvvm.read.ptx.sreg.tid.x does not have the result range attribute yet.

Full diff: https://github.com/llvm/llvm-project/pull/136297.diff

1 Files Affected:

  • (modified) clang/test/Headers/gpuintrin_lang.c (+1-1)
diff --git a/clang/test/Headers/gpuintrin_lang.c b/clang/test/Headers/gpuintrin_lang.c
index ab660ac5c8a49..b804d46071507 100644
--- a/clang/test/Headers/gpuintrin_lang.c
+++ b/clang/test/Headers/gpuintrin_lang.c
@@ -36,7 +36,7 @@ __device__ int foo() { return __gpu_thread_id_x(); }
 // CUDA-LABEL: define dso_local i32 @foo(
 // CUDA-SAME: ) #[[ATTR0:[0-9]+]] {
 // CUDA-NEXT:  [[ENTRY:.*:]]
-// CUDA-NEXT:    [[TMP0:%.*]] = call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
+// CUDA-NEXT:    [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
 // CUDA-NEXT:    ret i32 [[TMP0]]
 //
 // HIP-LABEL: define dso_local i32 @foo(

@jurahul jurahul merged commit e1b14d4 into llvm:main Apr 18, 2025
10 of 15 checks passed
@jurahul jurahul deleted the fix_nvptx_range_check branch April 18, 2025 12:01
@llvm-ci
Copy link
Collaborator

llvm-ci commented Apr 18, 2025

LLVM Buildbot has detected a new failure on builder premerge-monolithic-linux running on premerge-linux-1 while building clang at step 7 "test-build-unified-tree-check-all".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/153/builds/29242

Here is the relevant piece of the build log for the reference
Step 7 (test-build-unified-tree-check-all) failure: test (failure)
******************** TEST 'Clang :: Headers/gpuintrin_lang.c' FAILED ********************
Exit Code: 1

Command Output (stderr):
--
/build/buildbot/premerge-monolithic-linux/build/bin/clang -cc1 -internal-isystem /build/buildbot/premerge-monolithic-linux/build/lib/clang/21/include -nostdsysteminc -internal-isystem /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/Inputs/include    -internal-isystem /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/../../lib/Headers/cuda_wrappers    -internal-isystem /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/../../lib/Headers/    -fcuda-is-device -triple nvptx64 -emit-llvm /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/gpuintrin_lang.c -o -  | /build/buildbot/premerge-monolithic-linux/build/bin/FileCheck /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/gpuintrin_lang.c --check-prefix=CUDA # RUN: at line 2
+ /build/buildbot/premerge-monolithic-linux/build/bin/clang -cc1 -internal-isystem /build/buildbot/premerge-monolithic-linux/build/lib/clang/21/include -nostdsysteminc -internal-isystem /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/Inputs/include -internal-isystem /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/../../lib/Headers/cuda_wrappers -internal-isystem /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/../../lib/Headers/ -fcuda-is-device -triple nvptx64 -emit-llvm /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/gpuintrin_lang.c -o -
+ /build/buildbot/premerge-monolithic-linux/build/bin/FileCheck /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/gpuintrin_lang.c --check-prefix=CUDA
/build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/gpuintrin_lang.c:39:15: error: CUDA-NEXT: expected string not found in input
// CUDA-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
              ^
<stdin>:8:7: note: scanning from here
entry:
      ^
<stdin>:9:17: note: possible intended match here
 %0 = call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
                ^

Input file: <stdin>
Check file: /build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/gpuintrin_lang.c

-dump-input=help explains the following input dump.

Input was:
<<<<<<
           1: ; ModuleID = '/build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/gpuintrin_lang.c' 
           2: source_filename = "/build/buildbot/premerge-monolithic-linux/llvm-project/clang/test/Headers/gpuintrin_lang.c" 
           3: target datalayout = "e-p6:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64" 
           4: target triple = "nvptx64" 
           5:  
           6: ; Function Attrs: convergent noinline nounwind optnone 
           7: define dso_local i32 @foo() #0 { 
           8: entry: 
next:39'0           X error: no match found
           9:  %0 = call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() 
next:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
next:39'1                     ?                                                   possible intended match
          10:  ret i32 %0 
next:39'0     ~~~~~~~~~~~~
          11: } 
next:39'0     ~~
          12:  
next:39'0     ~
          13: ; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none) 
next:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
          14: declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1 
next:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           .
           .
           .
...

@llvm-ci
Copy link
Collaborator

llvm-ci commented Apr 18, 2025

LLVM Buildbot has detected a new failure on builder llvm-clang-x86_64-gcc-ubuntu running on sie-linux-worker3 while building clang at step 6 "test-build-unified-tree-check-all".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/174/builds/16421

Here is the relevant piece of the build log for the reference
Step 6 (test-build-unified-tree-check-all) failure: test (failure)
******************** TEST 'Clang :: Headers/gpuintrin_lang.c' FAILED ********************
Exit Code: 1

Command Output (stderr):
--
/home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/bin/clang -cc1 -internal-isystem /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/lib/clang/21/include -nostdsysteminc -internal-isystem /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/Inputs/include    -internal-isystem /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/../../lib/Headers/cuda_wrappers    -internal-isystem /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/../../lib/Headers/    -fcuda-is-device -triple nvptx64 -emit-llvm /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/gpuintrin_lang.c -o -  | /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/bin/FileCheck /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/gpuintrin_lang.c --check-prefix=CUDA # RUN: at line 2
+ /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/bin/clang -cc1 -internal-isystem /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/lib/clang/21/include -nostdsysteminc -internal-isystem /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/Inputs/include -internal-isystem /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/../../lib/Headers/cuda_wrappers -internal-isystem /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/../../lib/Headers/ -fcuda-is-device -triple nvptx64 -emit-llvm /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/gpuintrin_lang.c -o -
+ /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/build/bin/FileCheck /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/gpuintrin_lang.c --check-prefix=CUDA
�[1m/home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/gpuintrin_lang.c:39:15: �[0m�[0;1;31merror: �[0m�[1mCUDA-NEXT: expected string not found in input
�[0m// CUDA-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
�[0;1;32m              ^
�[0m�[1m<stdin>:8:7: �[0m�[0;1;30mnote: �[0m�[1mscanning from here
�[0mentry:
�[0;1;32m      ^
�[0m�[1m<stdin>:9:17: �[0m�[0;1;30mnote: �[0m�[1mpossible intended match here
�[0m %0 = call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
�[0;1;32m                ^
�[0m
Input file: <stdin>
Check file: /home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/gpuintrin_lang.c

-dump-input=help explains the following input dump.

Input was:
<<<<<<
�[1m�[0m�[0;1;30m            1: �[0m�[1m�[0;1;46m; ModuleID = '/home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/gpuintrin_lang.c' �[0m
�[0;1;30m            2: �[0m�[1m�[0;1;46msource_filename = "/home/buildbot/buildbot-root/llvm-clang-x86_64-gcc-ubuntu/llvm-project/clang/test/Headers/gpuintrin_lang.c" �[0m
�[0;1;30m            3: �[0m�[1m�[0;1;46mtarget datalayout = "e-p6:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64" �[0m
�[0;1;30m            4: �[0m�[1m�[0;1;46mtarget triple = "nvptx64" �[0m
�[0;1;30m            5: �[0m�[1m�[0;1;46m �[0m
�[0;1;30m            6: �[0m�[1m�[0;1;46m; Function Attrs: convergent noinline nounwind optnone �[0m
�[0;1;30m            7: �[0m�[1m�[0;1;46m�[0mdefine dso_local i32 @foo() #0 {�[0;1;46m �[0m
�[0;1;32mlabel:36'0     ^~~~~~~~~~~~~~~~~~~~~~~~~~
�[0m�[0;1;32mlabel:36'1     ^~~~~~~~~~~~~~~~~~~~~~~~~~
�[0m�[0;1;32msame:37'0                                ^~~~~~
�[0m�[0;1;32msame:37'1                                   ^    captured var "ATTR0"
�[0m�[0;1;30m            8: �[0m�[1m�[0;1;46m�[0mentry:�[0;1;46m �[0m
�[0;1;32mnext:38'0      ^~~~~~
�[0m�[0;1;32mnext:38'1      ^~~~~~  captured var "ENTRY"
�[0m�[0;1;31mnext:39'0            X error: no match found
�[0m�[0;1;30m            9: �[0m�[1m�[0;1;46m %0 = call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() �[0m
�[0;1;31mnext:39'0      ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
�[0m�[0;1;35mnext:39'1                      ?                                                   possible intended match
�[0m�[0;1;30m           10: �[0m�[1m�[0;1;46m ret i32 %0 �[0m
�[0;1;31mnext:39'0      ~~~~~~~~~~~~
�[0m�[0;1;30m           11: �[0m�[1m�[0;1;46m} �[0m
�[0;1;31mnext:39'0      ~~
�[0m�[0;1;30m           12: �[0m�[1m�[0;1;46m �[0m
�[0;1;31mnext:39'0      ~
�[0m�[0;1;30m           13: �[0m�[1m�[0;1;46m; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none) �[0m
...

@llvm-ci
Copy link
Collaborator

llvm-ci commented Apr 18, 2025

LLVM Buildbot has detected a new failure on builder clang-x86_64-debian-fast running on gribozavr4 while building clang at step 6 "test-build-unified-tree-check-all".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/56/builds/23794

Here is the relevant piece of the build log for the reference
Step 6 (test-build-unified-tree-check-all) failure: test (failure)
******************** TEST 'Clang :: Headers/gpuintrin_lang.c' FAILED ********************
Exit Code: 1

Command Output (stderr):
--
/b/1/clang-x86_64-debian-fast/llvm.obj/bin/clang -cc1 -internal-isystem /b/1/clang-x86_64-debian-fast/llvm.obj/lib/clang/21/include -nostdsysteminc -internal-isystem /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/Inputs/include    -internal-isystem /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/../../lib/Headers/cuda_wrappers    -internal-isystem /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/../../lib/Headers/    -fcuda-is-device -triple nvptx64 -emit-llvm /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/gpuintrin_lang.c -o -  | /b/1/clang-x86_64-debian-fast/llvm.obj/bin/FileCheck /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/gpuintrin_lang.c --check-prefix=CUDA # RUN: at line 2
+ /b/1/clang-x86_64-debian-fast/llvm.obj/bin/FileCheck /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/gpuintrin_lang.c --check-prefix=CUDA
+ /b/1/clang-x86_64-debian-fast/llvm.obj/bin/clang -cc1 -internal-isystem /b/1/clang-x86_64-debian-fast/llvm.obj/lib/clang/21/include -nostdsysteminc -internal-isystem /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/Inputs/include -internal-isystem /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/../../lib/Headers/cuda_wrappers -internal-isystem /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/../../lib/Headers/ -fcuda-is-device -triple nvptx64 -emit-llvm /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/gpuintrin_lang.c -o -
/b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/gpuintrin_lang.c:39:15: error: CUDA-NEXT: expected string not found in input
// CUDA-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
              ^
<stdin>:8:7: note: scanning from here
entry:
      ^
<stdin>:9:17: note: possible intended match here
 %0 = call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
                ^

Input file: <stdin>
Check file: /b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/gpuintrin_lang.c

-dump-input=help explains the following input dump.

Input was:
<<<<<<
           1: ; ModuleID = '/b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/gpuintrin_lang.c' 
           2: source_filename = "/b/1/clang-x86_64-debian-fast/llvm.src/clang/test/Headers/gpuintrin_lang.c" 
           3: target datalayout = "e-p6:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64" 
           4: target triple = "nvptx64" 
           5:  
           6: ; Function Attrs: convergent noinline nounwind optnone 
           7: define dso_local i32 @foo() #0 { 
           8: entry: 
next:39'0           X error: no match found
           9:  %0 = call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() 
next:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
next:39'1                     ?                                                   possible intended match
          10:  ret i32 %0 
next:39'0     ~~~~~~~~~~~~
          11: } 
next:39'0     ~~
          12:  
next:39'0     ~
          13: ; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none) 
next:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
          14: declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1 
next:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           .
           .
           .
...

@llvm-ci
Copy link
Collaborator

llvm-ci commented Apr 18, 2025

LLVM Buildbot has detected a new failure on builder llvm-x86_64-debian-dylib running on gribozavr4 while building clang at step 6 "test-build-unified-tree-check-clang".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/60/builds/25103

Here is the relevant piece of the build log for the reference
Step 6 (test-build-unified-tree-check-clang) failure: test (failure)
******************** TEST 'Clang :: Headers/gpuintrin_lang.c' FAILED ********************
Exit Code: 1

Command Output (stderr):
--
/b/1/llvm-x86_64-debian-dylib/build/bin/clang -cc1 -internal-isystem /b/1/llvm-x86_64-debian-dylib/build/lib/clang/21/include -nostdsysteminc -internal-isystem /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/Inputs/include    -internal-isystem /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/../../lib/Headers/cuda_wrappers    -internal-isystem /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/../../lib/Headers/    -fcuda-is-device -triple nvptx64 -emit-llvm /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/gpuintrin_lang.c -o -  | /b/1/llvm-x86_64-debian-dylib/build/bin/FileCheck /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/gpuintrin_lang.c --check-prefix=CUDA # RUN: at line 2
+ /b/1/llvm-x86_64-debian-dylib/build/bin/clang -cc1 -internal-isystem /b/1/llvm-x86_64-debian-dylib/build/lib/clang/21/include -nostdsysteminc -internal-isystem /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/Inputs/include -internal-isystem /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/../../lib/Headers/cuda_wrappers -internal-isystem /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/../../lib/Headers/ -fcuda-is-device -triple nvptx64 -emit-llvm /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/gpuintrin_lang.c -o -
+ /b/1/llvm-x86_64-debian-dylib/build/bin/FileCheck /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/gpuintrin_lang.c --check-prefix=CUDA
/b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/gpuintrin_lang.c:39:15: error: CUDA-NEXT: expected string not found in input
// CUDA-NEXT: [[TMP0:%.*]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
              ^
<stdin>:8:7: note: scanning from here
entry:
      ^
<stdin>:9:17: note: possible intended match here
 %0 = call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x()
                ^

Input file: <stdin>
Check file: /b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/gpuintrin_lang.c

-dump-input=help explains the following input dump.

Input was:
<<<<<<
           1: ; ModuleID = '/b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/gpuintrin_lang.c' 
           2: source_filename = "/b/1/llvm-x86_64-debian-dylib/llvm-project/clang/test/Headers/gpuintrin_lang.c" 
           3: target datalayout = "e-p6:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64" 
           4: target triple = "nvptx64" 
           5:  
           6: ; Function Attrs: convergent noinline nounwind optnone 
           7: define dso_local i32 @foo() #0 { 
           8: entry: 
next:39'0           X error: no match found
           9:  %0 = call range(i32 0, 1024) i32 @llvm.nvvm.read.ptx.sreg.tid.x() 
next:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
next:39'1                     ?                                                   possible intended match
          10:  ret i32 %0 
next:39'0     ~~~~~~~~~~~~
          11: } 
next:39'0     ~~
          12:  
next:39'0     ~
          13: ; Function Attrs: nocallback nofree nosync nounwind speculatable willreturn memory(none) 
next:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
          14: declare noundef i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1 
next:39'0     ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~
           .
           .
           .
...

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
clang Clang issues not falling into any other category
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants