n000l

Untitled

Jul 30th, 2021
46
Never
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
  1. #ifndef TMP_H
  2. #define TMP_H
  3.  
  4. #include "optixPathTracer.h"
  5. #include <cuda_runtime.h>
  6. #include <optix.h>
  7.  
  8. extern "C" {
  9. __constant__ Params params;
  10. }
  11.  
  12. __forceinline__ __device__ float other_function(
  13. const s0 *s0i1,
  14. const s0 *s0i0)
  15. {
  16. const s0 *curr = s0i0;
  17. const s0 *next = s0i1;
  18.  
  19. // loop runs for 2 iterations
  20. // in the first iteration, curr = s0i0 and next = s0i1
  21. // in the second iteration, curr = s0i1 and next = s0i1
  22. // note that the loop does not modify the s0 structs
  23. for (int index = 0; index < 2; index++) {
  24. // this assertion fails in the second iteration of the loop which
  25. // it should not. in the second iteration, current = s0i1
  26. // v1 should be 1.f, v2 should be false, v3 should be true
  27. GPU_ASSERT(curr->v1 != 0.f || (curr->v3 && curr->v2));
  28.  
  29. curr = next;
  30. }
  31.  
  32. return 0.f;
  33. }
  34.  
  35. __forceinline__ __device__ void eval()
  36. {
  37. s0 s0i0;
  38. s0i0.v1 = 1.f;
  39. s0i0.v2 = false;
  40. s0i0.v3 = false;
  41.  
  42. s3 *s3i0 = params.d_s3is;
  43. s0 s0i1 = s3i0->v6;
  44.  
  45. // s0i1 = params.d_s3is->v6 which was set on host side:
  46. // params.d_s3is->v6.v1 = undef
  47. // params.d_s3is->v6.v2 = false
  48. // params.d_s3is->v6.v3 = true
  49.  
  50. // hence the operand of the conditional operator below evals to false
  51. // hence s0i1.v1 is set to: 1.f * fabsf(-1.f) = 1.f
  52. s0i1.v1 = (s3i0->v6.v3 && s3i0->v6.v2) ? 0.f : 1.f * fabsf(-1.f);
  53.  
  54. // un-commenting this assertion makes the assertion in the
  55. // other_function succeed.
  56. // GPU_ASSERT(!(s3i0->v6.v3 && s3i0->v6.v2));
  57.  
  58. float ret = other_function(&s0i1, &s0i0);
  59. }
  60.  
  61. #endif
RAW Paste Data