Startseite Erkunden Hilfe
Anmelden
O3DE
/
DirectXShaderCompiler
Mirror von https://github.com/o3de/DirectXShaderCompiler
2
0
Fork 0
Dateien Issues 0 Wiki
Branch: version-2005
Branches Tags
DirectXShaderCo...
/
test
/
CodeGen
/
NVPTX
/
surf-read-cuda.ll

surf-read-cuda.ll 2.0 KB
Permalink Verlauf Originalformat

1234567891011121314151617181920212223242526272829303132333435363738394041424344454647484950515253 
  1. ; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s --check-prefix=SM20
    2. 

  2. ; RUN: llc < %s -march=nvptx -mcpu=sm_30 | FileCheck %s --check-prefix=SM30
    3. 

  3. 

  4. target triple = "nvptx-unknown-cuda"
    5. 

  5. 

  6. declare i32 @llvm.nvvm.suld.1d.i32.trap(i64, i32)
    7. 

  7. declare i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)*)
    8. 

  8. 

  9. 

  10. ; SM20-LABEL: .entry foo
    11. 

  11. ; SM30-LABEL: .entry foo
    12. 

  12. define void @foo(i64 %img, float* %red, i32 %idx) {
    13. 

  13. ; SM20: ld.param.u64    %rd[[SURFREG:[0-9]+]], [foo_param_0];
    14. 

  14. ; SM20: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFREG]], {%r{{[0-9]+}}}]
    15. 

  15. ; SM30: ld.param.u64    %rd[[SURFREG:[0-9]+]], [foo_param_0];
    16. 

  16. ; SM30: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFREG]], {%r{{[0-9]+}}}]
    17. 

  17.   %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx)
    18. 

  18. ; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
    19. 

  19. ; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
    20. 

  20.   %ret = sitofp i32 %val to float
    21. 

  21. ; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]]
    22. 

  22. ; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]]
    23. 

  23.   store float %ret, float* %red
    24. 

  24.   ret void
    25. 

  25. }
    26. 

  26. 

  27. @surf0 = internal addrspace(1) global i64 0, align 8
    28. 

  28. 

  29. ; SM20-LABEL: .entry bar
    30. 

  30. ; SM30-LABEL: .entry bar
    31. 

  31. define void @bar(float* %red, i32 %idx) {
    32. 

  32. ; SM30: mov.u64 %rd[[SURFHANDLE:[0-9]+]], surf0
    33. 

  33.   %surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @surf0)
    34. 

  34. ; SM20: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [surf0, {%r{{[0-9]+}}}]
    35. 

  35. ; SM30: suld.b.1d.b32.trap {%r[[RED:[0-9]+]]}, [%rd[[SURFHANDLE]], {%r{{[0-9]+}}}]
    36. 

  36.   %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %surfHandle, i32 %idx)
    37. 

  37. ; SM20: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
    38. 

  38. ; SM30: cvt.rn.f32.s32 %f[[REDF:[0-9]+]], %r[[RED]]
    39. 

  39.   %ret = sitofp i32 %val to float
    40. 

  40. ; SM20: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]]
    41. 

  41. ; SM30: st.global.f32 [%r{{[0-9]+}}], %f[[REDF]]
    42. 

  42.   store float %ret, float* %red
    43. 

  43.   ret void
    44. 

  44. }
    45. 

  45. 

  46. 

  47. 

  48. 

  49. !nvvm.annotations = !{!1, !2, !3}
    50. 

  50. !1 = !{void (i64, float*, i32)* @foo, !"kernel", i32 1}
    51. 

  51. !2 = !{void (float*, i32)* @bar, !"kernel", i32 1}
    52. 

  52. !3 = !{i64 addrspace(1)* @surf0, !"surface", i32 1}
    53. 

  53.