关于摩尔线程s80显卡OpenCL驱动崩溃的问题

2501_92453550 2025-06-14 15:04:34

本人长期开发、维护OpenCL开源可导光线追踪渲染器项目LibreDR https://codeberg.org/ybh1998/LibreDR 。最近在尝试适配摩尔线程s80显卡,在迁移摩尔线程s80时遇到编译kernel时驱动崩溃问题,详情如下:

本人在Ubuntu 20.04下分别尝试了musa_2.5.0和musa_2.7.1-rc3版本驱动,将kernel源代码使用clCreateProgramWithSource传入OpenCL驱动,kernel代码在附件“debug_kernel.cl”中。调用clBuildProgram编译kernel时程序发生segmentation fault,backtrace如附件“backtrace.png”中所示,显示segmentation fault发生在OpenCL驱动libufwriter_MUSA.so中,但由于缺失debug symbol,不能确定具体位置。

本人发现在clinfo中,摩尔线程显卡汇报支持cl_khr_spir扩展。本人尝试严格按照 https://github.com/KhronosGroup/SPIR 描述的标准编译SPIR程序,使用clCreateProgramWithBinary传入OpenCL驱动,在clBuildProgram时返回CL_DEVICE_NOT_AVAILABLE错误信息。

1. 使用源代码传入kernel时,由于segmentation fault发生在驱动代码内,是否方便处理相关驱动代码的问题?如果方便提供libufwriter_MUSA.so相关代码的debug symbol,本人可以协助调试;
2. 另外想问摩尔线程s80显卡OpenCL是否支持SPIR输入程序?如果支持,能否提供相应的使用样例?

 

 

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable

typedef struct {
  float3 r, rd, rn, render, render_weight;
  float3 render_last, render_last_weight;
  float2 rt;
  float pdf;
  ulong2 pcg32_state;
  int ri, hit_cnt;
} Ray;

typedef struct __attribute__((__packed__)) {
  const float vertex[8][3][170u];
  const float uv_xyz[4][256u][256u];
  const int bvh_left[86u], bvh_right[86u], bvh_next[86u];
  const float bvh_corner[3][2][86u];
} GeometryBuffer;

typedef struct __attribute__((__packed__)) {
  float ray[27u][262144u];
  const float texture[14][256u][256u];
  const uint texture_cdf[256u][256u];
    float d_texture[14][256u][256u][8u];
  const float envmap[3][6][512u][512u];
  const uint envmap_cdf[6][512u][512u];
  float d_envmap[3][6][512u][512u][8u];
} RenderBuffer;

typedef struct __attribute__((__packed__)) {
  float ray_r[3][2u][262144u][128u];
  uint ray_ri[2u][262144u][128u];
  float ray_pdf[2u][262144u][128u];
  float ray_render_last[3][262144u][128u];
} IntermediateBuffer;

uint pcg32_next(Ray *ray) {
  ulong old_state = ray->pcg32_state.x;
  ray->pcg32_state.x = old_state * 0x5851f42d4c957f2dul + ray->pcg32_state.y;
  uint xor_shifted = ((old_state >> 18u) ^ old_state) >> 27u;
  uint rot = old_state >> 59u;
  return (xor_shifted >> rot) | (xor_shifted << ((-rot) & 31));
}

float pcg32_next_float(Ray *ray) {
  union {
    uint u;
    float f;
  } x;
  x.u = (pcg32_next(ray) >> 9) | 0x3f800000u;
  return x.f - 1.f;
}

void pcg32_init(Ray *ray, ulong seed, ulong inc) {
  ray->pcg32_state.x = 0ul;
  ray->pcg32_state.y = (inc << 1u) | 1u;
  pcg32_next(ray);
  ray->pcg32_state.x += (0x853c49e6748fea9bul + seed);
  pcg32_next(ray);
}

float halton_float(const uint low_discrepancy, const uint base) {
  const uint PRIMES[] = {2, 3, 5, 7, 11, 13, 17, 19, 23, 29, 31, 37, 41, 43, 47, 53, 59, 61, 67, 71, 73, 79, 83, 89};
  const uint prime_base = PRIMES[base];
  float f = 1, r = 0;
  for (uint index = get_global_id(1) + low_discrepancy + 1; index > 0; index /= prime_base) {
    f /= prime_base;
    r += f * (index % prime_base);
  }
  return r;
}

uint halton(const uint low_discrepancy, const uint base) {
  return halton_float(low_discrepancy, base) * (float)UINT_MAX;
}

void atomic_add_f(volatile global float *source, const float operand) {
  if (operand == 0.f)
    return;
  union {
    uint u32;
    float f32;
  } next, expected, current;
  current.f32 = *source;
  do {
    expected.f32 = current.f32;
    next.f32 = expected.f32 + operand;
    current.u32 = atomic_cmpxchg((volatile __global uint *)source, expected.u32, next.u32);
  } while (current.u32 != expected.u32);
}

void local_gather_atomic_add_f (volatile global float *source, local float *operand_gather, const float operand) {
  operand_gather[get_local_id(1)] = operand;
  barrier(CLK_LOCAL_MEM_FENCE);
  for (int i = get_local_size(1) / 2; i > 0; i >>= 1) {
    if (get_local_id(1) < i)
      operand_gather[get_local_id(1)] += operand_gather[get_local_id(1) + i];
    barrier(CLK_LOCAL_MEM_FENCE);
  }
  if (get_local_id(1) == 0)
    atomic_add_f(source, operand_gather[0]);
  barrier(CLK_LOCAL_MEM_FENCE);
}

void d_material_normal(
    volatile global RenderBuffer *render_buffer,
    local float *operand_gather,
    const uint2 texture_index,
    const int i_bounce,
    const float3 d_normal) {
  const uint i_bucket = get_global_id(1) % 8u;
  atomic_add_f(&render_buffer->d_texture[0][texture_index.x][texture_index.y][i_bucket], d_normal.x);
  atomic_add_f(&render_buffer->d_texture[1][texture_index.x][texture_index.y][i_bucket], d_normal.y);
  atomic_add_f(&render_buffer->d_texture[2][texture_index.x][texture_index.y][i_bucket], d_normal.z);
}

float3 material_diffuse(
    const volatile global RenderBuffer *render_buffer,
    const uint2 texture_index,
    const int i_bounce) {
  return (float3)(render_buffer->texture[3][texture_index.x][texture_index.y],
                  render_buffer->texture[4][texture_index.x][texture_index.y],
                  render_buffer->texture[5][texture_index.x][texture_index.y]);
}

void d_material_diffuse(
    volatile global RenderBuffer *render_buffer,
    local float *operand_gather,
    const uint2 texture_index,
    const int i_bounce,
    const float3 d_diffuse) {
  const uint i_bucket = get_global_id(1) % 8u;
  atomic_add_f(&render_buffer->d_texture[3][texture_index.x][texture_index.y][i_bucket], d_diffuse.x);
  atomic_add_f(&render_buffer->d_texture[4][texture_index.x][texture_index.y][i_bucket], d_diffuse.y);
  atomic_add_f(&render_buffer->d_texture[5][texture_index.x][texture_index.y][i_bucket], d_diffuse.z);
}

float3 material_specular(
    const volatile global RenderBuffer *render_buffer,
    const uint2 texture_index,
    const int i_bounce) {
  return (float3)(render_buffer->texture[6][texture_index.x][texture_index.y],
                  render_buffer->texture[7][texture_index.x][texture_index.y],
                  render_buffer->texture[8][texture_index.x][texture_index.y]);
}

void d_material_specular(
    volatile global RenderBuffer *render_buffer,
    local float *operand_gather,
    const uint2 texture_index,
    const int i_bounce,
    const float3 d_specular) {
  const uint i_bucket = get_global_id(1) % 8u;
  atomic_add_f(&render_buffer->d_texture[6][texture_index.x][texture_index.y][i_bucket], d_specular.x);
  atomic_add_f(&render_buffer->d_texture[7][texture_index.x][texture_index.y][i_bucket], d_specular.y);
  atomic_add_f(&render_buffer->d_texture[8][texture_index.x][texture_index.y][i_bucket], d_specular.z);
}

float material_roughness(
    const volatile global RenderBuffer *render_buffer,
    const uint2 texture_index,
    const int i_bounce) {
  return render_buffer->texture[9][texture_index.x][texture_index.y];
}

void d_material_roughness(
    volatile global RenderBuffer *render_buffer,
    local float *operand_gather,
    const uint2 texture_index,
    const int i_bounce,
    const float d_roughness) {
  const uint i_bucket = get_global_id(1) % 8u;
  atomic_add_f(&render_buffer->d_texture[9][texture_index.x][texture_index.y][i_bucket], d_roughness);
}

float phong_roughness_to_shininess(const float roughness) {
  return 2.f / roughness - 2.f;
}

float d_phong_roughness_to_shininess(const float roughness, const float d_shininess) {
  return -2.f * d_shininess / pown(roughness, 2);
}

float specular_D_beckmann(const float3 rn, const float3 h, const float roughness) {
  const float alpha = roughness / 4.f;
  const float n_dot_h = dot(rn, h);
  return exp((pown(n_dot_h, 2) - 1.f) / (alpha * pown(n_dot_h, 2))) /
    (4 * M_PI_F * alpha * pown(n_dot_h, 4));
}

void d_specular_D_beckmann(
    const float3 rn,
    const float3 h,
    const float roughness,
    const float d_D,
    float3 *d_normal,
    float *d_roughness) {
  const float alpha = roughness / 4.f;
  const float n_dot_h = dot(rn, h);
  const float d_n_dot_h = d_D * exp((pown(n_dot_h, 2) - 1.f) / (alpha * pown(n_dot_h, 2))) *
    (1.f - 2.f * alpha * pown(n_dot_h, 2)) / (2.f * M_PI_F * pown(alpha, 2) * pown(n_dot_h, 7));
  *d_normal += d_n_dot_h * h;
  const float d_alpha = d_D * exp((pown(n_dot_h, 2) - 1.f) / (alpha * pown(n_dot_h, 2))) *
    (1.f - (alpha + 1.f) * pown(n_dot_h, 2)) / (4.f * M_PI_F * pown(n_dot_h, 6) * pow(alpha, 3));
  *d_roughness += d_alpha / 4.f;
}

int bin_search(const global uint *cdf, const uint size, const uint value) {
  int a = -1, b = size - 1;
  while (a + 1 < b) {
    int m = (a + b) / 2;
    if (cdf[m] < value)
      a = m;
    else
      b = m;
  }
  return b;
}

float3 envmap_offset_to_rd(const uint envmap_offset) {
  const uint envmap_face_id = envmap_offset / (512u * 512u);
  const uint envmap_row = (envmap_offset / 512u) % 512u;
  const uint envmap_col = envmap_offset % 512u;
  const float2 background_uv = (float2)(
    (envmap_col + 0.5f) / 512u,
    (512u - 0.5f - envmap_row) / 512u) * 2.f - 1.f;
  float3 next_rd;
  switch (envmap_face_id / 2) {
    case 0:
      next_rd.x = (envmap_face_id % 2) ? -1.f : 1.f;
      next_rd.y = background_uv.x;
      next_rd.z = background_uv.y;
      break;
    case 1:
      next_rd.x = background_uv.x;
      next_rd.y = (envmap_face_id % 2) ? -1.f : 1.f;
      next_rd.z = background_uv.y;
      break;
    case 2:
      next_rd.x = background_uv.x;
      next_rd.y = background_uv.y;
      next_rd.z = (envmap_face_id % 2) ? -1.f : 1.f;
      break;
  }
  next_rd /= sqrt(dot(next_rd, next_rd));
  return next_rd;
}

float3 sample_reflection(
    const global GeometryBuffer *geometry_buffer,
    const volatile global RenderBuffer *render_buffer,
    Ray *ray,
    const uint2 texture_index,
    int i_bounce,
    uint low_discrepancy) {
  const float3 r = ray->rd - 2 * dot(ray->rd, ray->rn) * ray->rn;
  const uint w = i_bounce < 0 ? halton(low_discrepancy, i_bounce * 3 + 2) : pcg32_next(ray);
  float tmp;
  float3 next_rd = (float3)(0.f, 0.f, 0.f);
  if (w < 1073741826u + 1073741823u) {
    float3 n_coord[3];
    n_coord[2] = ray->rn;
    n_coord[0] = cross((float3)(0.f, 0.f, 1.f), n_coord[2]);
    tmp = dot(n_coord[0], n_coord[0]);
    if (tmp < 1e-3f) {
      n_coord[0] = cross((float3)(0.f, 1.f, 0.f), n_coord[2]);
      tmp = dot(n_coord[0], n_coord[0]);
    }
    if (tmp < 1e-3f) {
      n_coord[0] = cross((float3)(1.f, 0.f, 0.f), n_coord[2]);
      tmp = dot(n_coord[0], n_coord[0]);
    }
    n_coord[0] /= sqrt(tmp);
    n_coord[1] = cross(n_coord[0], n_coord[2]);
    n_coord[1] /= sqrt(dot(n_coord[1], n_coord[1]));
    const float sample_a =
      i_bounce < 0 ? halton_float(low_discrepancy, i_bounce * 3 + 3) : pcg32_next_float(ray);
    const float sample_b =
      i_bounce < 0 ? halton_float(low_discrepancy, i_bounce * 3 + 4) : pcg32_next_float(ray);
    const float2 sample = (float2)(sample_a, sample_b);
    const float sin_phi = sin(2.f * M_PI_F * sample.x);
    const float cos_phi = cos(2.f * M_PI_F * sample.x);
    if (w <= 1073741826u) {
      tmp = sqrt(1.f - sample.y);
      const float3 local_dir = (float3)(cos_phi * tmp, sin_phi * tmp, sqrt(sample.y));
      next_rd = local_dir.x * n_coord[0] + local_dir.y * n_coord[1] + local_dir.z * n_coord[2];
    } else {
      float3 r_coord[3];
      r_coord[2] = r;
      r_coord[0] = cross((float3)(0.f, 0.f, 1.f), r_coord[2]);
      tmp = dot(r_coord[0], r_coord[0]);
      if (tmp < 1e-3f) {
        r_coord[0] = cross((float3)(0.f, 1.f, 0.f), r_coord[2]);
        tmp = dot(r_coord[0], r_coord[0]);
      }
      if (tmp < 1e-3f) {
        r_coord[0] = cross((float3)(1.f, 0.f, 0.f), r_coord[2]);
        tmp = dot(r_coord[0], r_coord[0]);
      }
      r_coord[0] /= sqrt(tmp);
      r_coord[1] = cross(r_coord[0], r_coord[2]);
      r_coord[1] /= sqrt(dot(r_coord[1], r_coord[1]));
      const float roughness = material_roughness(render_buffer, texture_index, i_bounce);
      const float shininess = phong_roughness_to_shininess(roughness);
      const float cos_theta = pow(sample.y, 1.f / (shininess + 1.f));
      const float sin_theta = sqrt(1.f - cos_theta * cos_theta);
      const float3 local_dir = (float3)(cos_phi * sin_theta, sin_phi * sin_theta, cos_theta);
      next_rd = local_dir.x * r_coord[0] + local_dir.y * r_coord[1] + local_dir.z * r_coord[2];
      if (dot(next_rd, ray->rn) < 0) {
        next_rd -= 2 * dot(next_rd, ray->rn) * ray->rn;
      }
    }
  } else if (w < 1073741826u + 1073741823u + 1073741823u) {
    const uint texture_offset = bin_search(
      (const global uint *)render_buffer->texture_cdf,
      256u * 256u,
      i_bounce < 0 ? halton(low_discrepancy, i_bounce * 3 + 3) : pcg32_next(ray));
    const uint2 texture_index = (uint2)(texture_offset / 256u, texture_offset % 256u);
    const float3 look_at = (float3)(
      geometry_buffer->uv_xyz[0][texture_index.x][texture_index.y],
      geometry_buffer->uv_xyz[1][texture_index.x][texture_index.y],
      geometry_buffer->uv_xyz[2][texture_index.x][texture_index.y]);
    next_rd = look_at - ray->r;
    const float dist = sqrt(dot(next_rd, next_rd));
    if (dist < 1e-3f) {
      ray->render_weight = (float3)(0.f, 0.f, 0.f);
      next_rd = (float3)(0.f, 1.f, 0.f);
    } else {
      next_rd /= dist;
    }
  } else {
    const uint envmap_offset = bin_search(
      (const global uint *)render_buffer->envmap_cdf,
      6 * 512u * 512u,
      i_bounce < 0 ? halton(low_discrepancy, i_bounce * 3 + 3) : pcg32_next(ray));
    next_rd = envmap_offset_to_rd(envmap_offset);
  }
  return next_rd;
}

float3 diffuse_brdf(
    const volatile global RenderBuffer *render_buffer,
    Ray *ray,
    const float3 next_rd,
    const uint2 texture_index,
    const int i_bounce) {
  float3 brdf = (float3)(0.f, 0.f, 0.f);
  if (dot(ray->rn, next_rd) < 1e-3f)
    return brdf;
  const float3 diffuse = material_diffuse(render_buffer, texture_index, i_bounce);
  brdf = diffuse * dot(ray->rn, next_rd) / M_PI_F;
  return brdf;
}

void d_diffuse_brdf(
    volatile global RenderBuffer *render_buffer,
    local float *operand_gather,
    Ray *ray,
    float3 next_rd,
    const uint2 texture_index,
    const int i_bounce,
    float3 d_brdf) {
  float3 d_diffuse = (float3)(0.f, 0.f, 0.f);
  float3 d_normal = (float3)(0.f, 0.f, 0.f);
  if (fabs(d_brdf.x) < 1e-6f && fabs(d_brdf.y) < 1e-6f && fabs(d_brdf.z) < 1e-6f)
    goto d_return;
  if (dot(ray->rn, next_rd) < 1e-3f)
    goto d_return;
  d_diffuse += d_brdf * dot(ray->rn, next_rd) / M_PI_F;
  const float3 diffuse = material_diffuse(render_buffer, texture_index, i_bounce);
  d_normal += next_rd * dot(d_brdf, diffuse) / M_PI_F;
d_return:
  d_material_diffuse(render_buffer, operand_gather, texture_index, i_bounce, d_diffuse);
  d_material_normal(render_buffer, operand_gather, texture_index, i_bounce, d_normal);
}

float3 specular_brdf(
    const volatile global RenderBuffer *render_buffer,
    Ray *ray,
    const float3 next_rd,
    const uint2 texture_index,
    const int i_bounce) {
  float3 brdf = (float3)(0.f, 0.f, 0.f);
  if (dot(ray->rn, next_rd) < 1e-3f)
    return brdf;
  const float3 specular = material_specular(render_buffer, texture_index, i_bounce);
  const float roughness = material_roughness(render_buffer, texture_index, i_bounce);
  // TODO: Shall we move the data EPS check to rust code?
  if (roughness < 1e-6f)
    return brdf;
  float3 h = next_rd - ray->rd;
  if (dot(h, h) < 1e-3f)
    return brdf;
  h /= sqrt(dot(h, h));
  if (dot(h, -ray->rd) < 1e-3f || dot(h, ray->rn) < 1e-3f)
    return brdf;
  const float3 F = specular + (1.f - specular) * pown(1.f - dot(ray->rn, -ray->rd), 5);
  const float G1 = 2.f * dot(h, ray->rn) * dot(ray->rn, -ray->rd) / dot(h, -ray->rd);
  const float G2 = 2.f * dot(h, ray->rn) * dot(ray->rn, next_rd) / dot(h, -ray->rd);
  const float G = min(1.f, min(G1, G2));
  const float D = specular_D_beckmann(ray->rn, h, roughness);
  brdf = D * G * F / dot(ray->rn, -ray->rd);
  return brdf;
}

void d_specular_brdf(
    volatile global RenderBuffer *render_buffer,
    local float *operand_gather,
    Ray *ray,
    float3 next_rd,
    const uint2 texture_index,
    const int i_bounce,
    float3 d_brdf) {
  float3 d_specular = (float3)(0.f, 0.f, 0.f);
  float d_roughness = 0.f;
  float3 d_normal = (float3)(0.f, 0.f, 0.f);
  if (fabs(d_brdf.x) < 1e-6f && fabs(d_brdf.y) < 1e-6f && fabs(d_brdf.z) < 1e-6f)
    goto d_return;
  if (dot(ray->rn, next_rd) < 1e-3f)
    goto d_return;
  const float3 specular = material_specular(render_buffer, texture_index, i_bounce);
  float roughness = material_roughness(render_buffer, texture_index, i_bounce);
  // TODO: Shall we move the data EPS check to rust code?
  if (roughness < 1e-6f)
    goto d_return;
  float3 h = next_rd - ray->rd;
  if (dot(h, h) < 1e-3f)
    goto d_return;
  h /= sqrt(dot(h, h));
  if (dot(h, -ray->rd) < 1e-3f || dot(h, ray->rn) < 1e-3f)
    goto d_return;
  const float3 F = specular + (1.f - specular) * pown(1.f - dot(ray->rn, -ray->rd), 5);
  const float G1 = 2.f * dot(h, ray->rn) * dot(ray->rn, -ray->rd) / dot(h, -ray->rd);
  const float G2 = 2.f * dot(h, ray->rn) * dot(ray->rn, next_rd) / dot(h, -ray->rd);
  const float G = min(1.f, min(G1, G2));
  const float D = specular_D_beckmann(ray->rn, h, roughness);
  const float d_D = dot(d_brdf, F) * G / dot(ray->rn, -ray->rd);
  float3 d_n_dot_v = -dot(d_brdf, F) * D * G / pown(dot(ray->rn, -ray->rd), 2);
  const float3 d_F = d_brdf * D * G / dot(ray->rn, -ray->rd);
  d_specular = d_F * (1.f - pown(1.f - dot(ray->rn, -ray->rd), 5));
  d_n_dot_v += -5 * dot(d_F, 1.f - specular) * pown(1.f - dot(ray->rn, -ray->rd), 4);
  d_normal += d_n_dot_v * -ray->rd;
  const float d_G = dot(d_brdf, F) * D / dot(ray->rn, -ray->rd);
  if (G1 < 1.f && G1 <= G2) {
    d_normal += 2.f * d_G * (dot(h, ray->rn) * -ray->rd + h * dot(ray->rn, -ray->rd)) / dot(h, -ray->rd);
  }
  if (G2 < 1.f && G2 < G1) {
    d_normal += 2.f * d_G * (dot(h, ray->rn) * next_rd + h * dot(ray->rn, next_rd)) / dot(h, -ray->rd);
  }
  d_specular_D_beckmann(ray->rn, h, roughness, d_D, &d_normal, &d_roughness);
d_return:
  d_material_specular(render_buffer, operand_gather, texture_index, i_bounce, d_specular);
  d_material_roughness(render_buffer, operand_gather, texture_index, i_bounce, d_roughness);
  d_material_normal(render_buffer, operand_gather, texture_index, i_bounce, d_normal);
}

float diffuse_pdf(Ray *ray, const float3 next_rd) {
  const float pdf = max(dot(ray->rn, next_rd), 0.f) / M_PI_F;
  return 0.25000000052386895f * pdf;
}

float specular_pdf(
    const volatile global RenderBuffer *render_buffer,
    Ray *ray,
    const float3 next_rd,
    const uint2 texture_index,
    const int i_bounce) {
  float pdf = 0.f;
  if (dot(ray->rn, next_rd) < 1e-3f)
    return pdf;
  const float roughness = material_roughness(render_buffer, texture_index, i_bounce);
  const float shininess = phong_roughness_to_shininess(roughness);
  const float3 r = ray->rd - 2 * dot(ray->rd, ray->rn) * ray->rn;
  float r_dot_wi = max(dot(r, next_rd - 2 * dot(next_rd, ray->rn) * ray->rn), 0.f);
  pdf += pow(r_dot_wi, shininess) * (shininess + 1.f) / (2.f * M_PI_F);
  r_dot_wi = max(dot(r, next_rd), 0.f);
  pdf += pow(r_dot_wi, shininess) * (shininess + 1.f) / (2.f * M_PI_F);
  return 0.24999999982537702f * pdf;
}

float float3_access(const float3 data, const uint idx) {
  if (idx == 0)
    return data.x;
  if (idx == 1)
    return data.y;
  return data.z;
}

float envmap_pdf(const global uint *cdf, const Ray *ray, const uint3 envmap_index) {
  uint envmap_offset = envmap_index.x * 512u * 512u +
    envmap_index.y * 512u + envmap_index.z;
  uint pdf = cdf[envmap_offset];
  if (envmap_offset > 0)
    pdf -= cdf[envmap_offset - 1];
  float n_dot_v = max(fabs(ray->rd.x), max(fabs(ray->rd.y), fabs(ray->rd.z)));
  return 3.814697263848643e-6f * pdf / pown(n_dot_v, 3);
}

void background_intersect(const volatile global RenderBuffer *render_buffer, Ray *ray, uint3 *envmap_index) {
  uint face_id = 0;
  const float3 rd_abs = fabs(ray->rd);
  if (rd_abs.x > rd_abs.y) {
    if (rd_abs.x > rd_abs.z)
      face_id = 0;
    else
      face_id = 2;
  } else {
    if (rd_abs.y > rd_abs.z)
      face_id = 1;
    else
      face_id = 2;
  }
  const float3 rd = ray->rd / float3_access(rd_abs, face_id);
  float2 uv;
  switch (face_id) {
  case 0:
    uv = (float2)(rd.y, rd.z);
    break;
  case 1:
    uv = (float2)(rd.x, rd.z);
    break;
  case 2:
    uv = (float2)(rd.x, rd.y);
    break;
  }
  uv = (uv + 1.f) / 2.f;
  if (float3_access(ray->rd, face_id) < 0)
    face_id = face_id * 2 + 1;
  else
    face_id = face_id * 2;
  uv = clamp(uv, 1e-6f, 1.f - 1e-6f);
  envmap_index->x = face_id;
  envmap_index->y = 512u - 1 - floor(uv.y * 512u);
  envmap_index->z = floor(uv.x * 512u);
  ray->pdf += envmap_pdf((const global uint *)render_buffer->envmap_cdf, ray, *envmap_index);
}

float triangle_intersect(const global GeometryBuffer *geometry_buffer, Ray *ray, const int triangle_id) {
  const float3 tri_v0 = (float3)(
    geometry_buffer->vertex[0][0][triangle_id],
    geometry_buffer->vertex[1][0][triangle_id],
    geometry_buffer->vertex[2][0][triangle_id]);
  const float3 u = (float3)(
    geometry_buffer->vertex[0][1][triangle_id],
    geometry_buffer->vertex[1][1][triangle_id],
    geometry_buffer->vertex[2][1][triangle_id]) - tri_v0;
  const float3 v = (float3)(
    geometry_buffer->vertex[0][2][triangle_id],
    geometry_buffer->vertex[1][2][triangle_id],
    geometry_buffer->vertex[2][2][triangle_id]) - tri_v0;
  const float3 fn = cross(u, v);
  const float3 w0 = ray->r - tri_v0;
  const float a = -dot(fn, w0);
  const float b = dot(fn, ray->rd);
  if (fabs(b) < 1e-12f) return MAXFLOAT;
  const float r = a / b;
  if (r <= 0) return MAXFLOAT;
  const float3 I = ray->r + ray->rd * r;
  const float uu = dot(u, u);
  const float uv = dot(u, v);
  const float vv = dot(v, v);
  const float3 w = I - tri_v0;
  const float wu = dot(w, u);
  const float wv = dot(w, v);
  const float D = uu * vv - uv * uv;
  const float D_min = -1e-5f * D;
  const float D_max = D - D_min;
  const float s = vv * wu - uv * wv;
  if (s < D_min || s > D_max) return MAXFLOAT;
  const float t = uu * wv - uv * wu;
  if (t < D_min || (s + t) > D_max) return MAXFLOAT;
  return r;
}

bool aabb_intersect(const global GeometryBuffer *geometry_buffer, Ray *ray, const int bvh_id) {
  const float3 corner[2] = {
    (float3)(
      geometry_buffer->bvh_corner[0][0][bvh_id],
      geometry_buffer->bvh_corner[1][0][bvh_id],
      geometry_buffer->bvh_corner[2][0][bvh_id]),
    (float3)(
      geometry_buffer->bvh_corner[0][1][bvh_id],
      geometry_buffer->bvh_corner[1][1][bvh_id],
      geometry_buffer->bvh_corner[2][1][bvh_id])
  };
  const float3 dir_frac = 1.0f / ray->rd;
  const float3 tmp[2] = {
    (corner[0] - ray->r) * dir_frac,
    (corner[1] - ray->r) * dir_frac
  };
  const float t_min = max(max(min(tmp[0].x, tmp[1].x), min(tmp[0].y, tmp[1].y)), min(tmp[0].z, tmp[1].z));
  const float t_max = min(min(max(tmp[0].x, tmp[1].x), max(tmp[0].y, tmp[1].y)), max(tmp[0].z, tmp[1].z));
  return t_max >= t_min;
}

float update_ray(
    const global GeometryBuffer *geometry_buffer,
    Ray *ray,
    const float d,
    const int triangle_id,
    uint2 *texture_index) {
  const float3 next_r = ray->r + ray->rd * d;
  const float3 tri_v0 = (float3)(
    geometry_buffer->vertex[0][0][triangle_id],
    geometry_buffer->vertex[1][0][triangle_id],
    geometry_buffer->vertex[2][0][triangle_id]);
  const float3 tri_v1 = (float3)(
    geometry_buffer->vertex[0][1][triangle_id],
    geometry_buffer->vertex[1][1][triangle_id],
    geometry_buffer->vertex[2][1][triangle_id]);
  const float3 tri_v2 = (float3)(
    geometry_buffer->vertex[0][2][triangle_id],
    geometry_buffer->vertex[1][2][triangle_id],
    geometry_buffer->vertex[2][2][triangle_id]);
  const float2 tri_vt0 = (float2)(
    geometry_buffer->vertex[6][0][triangle_id],
    geometry_buffer->vertex[7][0][triangle_id]);
  const float2 tri_vt1 = (float2)(
    geometry_buffer->vertex[6][1][triangle_id],
    geometry_buffer->vertex[7][1][triangle_id]);
  const float2 tri_vt2 = (float2)(
    geometry_buffer->vertex[6][2][triangle_id],
    geometry_buffer->vertex[7][2][triangle_id]);
  const float3 u_cross_v = cross(tri_v1 - tri_v0, tri_v2 - tri_v0);
  const float area_v = sqrt(dot(u_cross_v, u_cross_v));
  const float area_vt = fabs((tri_vt1 - tri_vt0).x * (tri_vt2 - tri_vt0).y -
                             (tri_vt1 - tri_vt0).y * (tri_vt2 - tri_vt0).x);
  float w = dot(cross(tri_v1 - tri_v0, next_r - tri_v0), u_cross_v) / dot(u_cross_v, u_cross_v);
  float v = dot(cross(next_r - tri_v0, tri_v2 - tri_v0), u_cross_v) / dot(u_cross_v, u_cross_v);
  w = max(1e-6f, min(1.f - 1e-6f, w));
  v = max(1e-6f, min(1.f - 1e-6f - w, v));
  float u = 1.f - w - v;
  ray->rt = tri_vt0 * u + tri_vt1 * v + tri_vt2 * w;
  ray->rt = clamp(ray->rt, 1e-6f, 1.f - 1e-6f);
  texture_index->x = 256u - 1 - floor(ray->rt.y * 256u);
  texture_index->y = floor(ray->rt.x * 256u);
  ray->rn = u_cross_v / sqrt(dot(u_cross_v, u_cross_v));
  if (dot(ray->rd, ray->rn) > 0)
    ray->rn = -ray->rn;
  return area_vt / area_v;
}

float texture_pdf(
    const global uint *texture_cdf,
    const Ray *ray,
    const float d,
    const float texture_area_ratio,
    const uint2 texture_index) {
  const uint texture_offset = texture_index.x * 256u + texture_index.y;
  uint pdf = texture_cdf[texture_offset];
  if (texture_offset > 0)
    pdf -= texture_cdf[texture_offset - 1];
  float n_dot_v = fabs(dot(ray->rd, ray->rn));
  if (n_dot_v < 1e-3f)
    return MAXFLOAT;
  return 3.814697263848643e-6f * pdf * texture_area_ratio * pown(d, 2) / n_dot_v;
}

void any_hit(
    const global GeometryBuffer *geometry_buffer,
    const global uint *texture_cdf,
    Ray *ray,
    const float d,
    const int triangle_id,
    uint2 *texture_index) {
  const float texture_area_ratio = update_ray(geometry_buffer, ray, d, triangle_id, texture_index);
  ray->pdf += texture_pdf(texture_cdf, ray, d, texture_area_ratio, *texture_index);
}

void closest_hit(
    const global GeometryBuffer *geometry_buffer,
    Ray *ray,
    const float d,
    uint2 *texture_index) {
  update_ray(geometry_buffer, ray, d, ray->ri, texture_index);
  ray->r += ray->rd * d;
}

void miss(Ray *ray) {
  ray->rn = (float3)(0.0f, 0.0f, 0.0f);
  ray->rt = (float2)(0.0f, 0.0f);
  ray->ri = -1;
}

float scene_intersect(
    const global GeometryBuffer *geometry_buffer,
    const global uint *texture_cdf,
    Ray *ray,
    uint2 *texture_index) {
  if (ray->ri == -2)
    return 0.f;
  ray->hit_cnt = 0;
  float d = MAXFLOAT;
  int bvh_id = 0;
  while (bvh_id != -1) {
    int bvh_left = geometry_buffer->bvh_left[bvh_id];
    int bvh_right = geometry_buffer->bvh_right[bvh_id];
    int bvh_next = geometry_buffer->bvh_next[bvh_id];
    bool has_intersect = aabb_intersect(geometry_buffer, ray, bvh_id);
    bool bvh_is_leaf = (bvh_right != -1);
    if (!has_intersect) {
      bvh_id = bvh_next;
      continue;
    }
    if (!bvh_is_leaf) {
      bvh_id = bvh_left;
      continue;
    }
    for (int i = bvh_left; i < bvh_right; ++i) {
      float new_d = triangle_intersect(geometry_buffer, ray, i);
      if (new_d != MAXFLOAT) {
        if (new_d < d) {
          ray->ri = i;
          d = new_d;
        }
        ray->hit_cnt++;
        any_hit(geometry_buffer, texture_cdf, ray, new_d, i, texture_index);
      }
    }
    bvh_id = bvh_next;
  }
  if (ray->hit_cnt) {
    closest_hit(geometry_buffer, ray, d, texture_index);
    return d;
  }
  miss(ray);
  return 0.f;
}

Ray ray_init(const int srand, const uint low_discrepancy, const volatile global RenderBuffer *render_buffer) {
  Ray ray;
  const uint i_pixel = get_global_id(0);
  const uint i_sample = get_global_id(1);
  if (srand >= 0)
    pcg32_init(&ray, srand, i_sample);
  else
    pcg32_init(&ray, i_pixel + get_global_size(0) * (1 - srand), i_sample);
    float2 pixel_coord = (float2)(pcg32_next_float(&ray), pcg32_next_float(&ray));
  if (i_sample == 0) {
    pixel_coord = (float2)(0.5f, 0.5f);
  }
  ray.r = (float3)(render_buffer->ray[0][i_pixel],
                   render_buffer->ray[1][i_pixel],
                   render_buffer->ray[2][i_pixel]
  ) + pixel_coord.x * (float3)(render_buffer->ray[3][i_pixel],
                               render_buffer->ray[4][i_pixel],
                               render_buffer->ray[5][i_pixel]
  ) + pixel_coord.y * (float3)(render_buffer->ray[6][i_pixel],
                               render_buffer->ray[7][i_pixel],
                               render_buffer->ray[8][i_pixel]);
  ray.rd = (float3)(render_buffer->ray[9][i_pixel],
                    render_buffer->ray[10][i_pixel],
                    render_buffer->ray[11][i_pixel]
  ) + pixel_coord.x * (float3)(render_buffer->ray[12][i_pixel],
                               render_buffer->ray[13][i_pixel],
                               render_buffer->ray[14][i_pixel]
  ) + pixel_coord.y * (float3)(render_buffer->ray[15][i_pixel],
                               render_buffer->ray[16][i_pixel],
                               render_buffer->ray[17][i_pixel]);
  ray.rd /= sqrt(dot(ray.rd, ray.rd));
  ray.render = (float3)(0.f, 0.f, 0.f);
  ray.render_last = (float3)(0.f, 0.f, 0.f);
  ray.render_last_weight = (float3)(1.f, 1.f, 1.f);
  ray.ri = 0;
  return ray;
}

void material_normal(
    Ray *ray,
    const volatile global RenderBuffer *render_buffer,
    const uint2 texture_index,
    const int i_bounce) {
  if (!(false && i_bounce == 0) && ray->ri >= 0)
    ray->rn = (float3)(render_buffer->texture[0][texture_index.x][texture_index.y],
                       render_buffer->texture[1][texture_index.x][texture_index.y],
                       render_buffer->texture[2][texture_index.x][texture_index.y]);
  if (dot(ray->rn, ray->rn) < 1e-3f) {
    ray->ri = -2;
  } else {
    ray->rn /= sqrt(dot(ray->rn, ray->rn));
    if (dot(ray->rd, ray->rn) > 0)
      ray->rn = -ray->rn;
    if (dot(ray->rd, ray->rn) > -1e-3f)
      ray->ri = -2;
  }
}

float3 material_intensity(const volatile global RenderBuffer *render_buffer, const uint2 texture_index, const int i_bounce) {
  return (float3)(render_buffer->texture[10][texture_index.x][texture_index.y],
                  render_buffer->texture[11][texture_index.x][texture_index.y],
                  render_buffer->texture[12][texture_index.x][texture_index.y]);
}

float material_window(const volatile global RenderBuffer *render_buffer, const uint2 texture_index, const int i_bounce) {
  return render_buffer->texture[13][texture_index.x][texture_index.y];
}

float3 material_envmap(const volatile global RenderBuffer *render_buffer, const uint3 envmap_index) {
  return (float3)(render_buffer->envmap[0][envmap_index.x][envmap_index.y][envmap_index.z],
                  render_buffer->envmap[1][envmap_index.x][envmap_index.y][envmap_index.z],
                  render_buffer->envmap[2][envmap_index.x][envmap_index.y][envmap_index.z]);
}

void d_material_intensity(
      volatile global RenderBuffer *render_buffer,
      local float *operand_gather,
      const uint2 texture_index,
      const int i_bounce,
      const float3 d_intensity) {
  const uint i_bucket = get_global_id(1) % 8u;
  atomic_add_f(&render_buffer->d_texture[10][texture_index.x][texture_index.y][i_bucket], d_intensity.x);
  atomic_add_f(&render_buffer->d_texture[11][texture_index.x][texture_index.y][i_bucket], d_intensity.y);
  atomic_add_f(&render_buffer->d_texture[12][texture_index.x][texture_index.y][i_bucket], d_intensity.z);
}

void d_material_window(
      volatile global RenderBuffer *render_buffer,
      local float *operand_gather,
      const uint2 texture_index,
      const int i_bounce,
      const float d_window) {
  const uint i_bucket = get_global_id(1) % 8u;
  atomic_add_f(&render_buffer->d_texture[13][texture_index.x][texture_index.y][i_bucket], d_window);
}

void d_material_envmap(volatile global RenderBuffer *render_buffer, const uint3 envmap_index, const float3 d_envmap) {
  const uint i_bucket = get_global_id(1) % 8u;
  atomic_add_f(&render_buffer->d_envmap[0][envmap_index.x][envmap_index.y][envmap_index.z][i_bucket], d_envmap.x);
  atomic_add_f(&render_buffer->d_envmap[1][envmap_index.x][envmap_index.y][envmap_index.z][i_bucket], d_envmap.y);
  atomic_add_f(&render_buffer->d_envmap[2][envmap_index.x][envmap_index.y][envmap_index.z][i_bucket], d_envmap.z);
}

kernel void kernel_ray_tracing_forward(
    const int srand,
    const uint low_discrepancy,
    const global GeometryBuffer *restrict geometry_buffer,
    volatile global RenderBuffer *restrict render_buffer,
    global IntermediateBuffer *restrict intermediate_buffer) {
  local float operand_gather[128u];
  const uint i_pixel = get_global_id(0);
  const uint i_sample = get_global_id(1);
  Ray ray = ray_init(srand, low_discrepancy, render_buffer);
  for (int i_bounce = 0; i_bounce <= 4u; ++i_bounce) {
    const float clip_near = i_bounce == 0 ? 0.01f : i_bounce == 1 ? 0.0001f : 0.0001f;
    float d;
    uint2 texture_index;
    uint3 envmap_index;
    background_intersect(render_buffer, &ray, &envmap_index);
    if (false && i_bounce == 0) {
      d = render_buffer->ray[23][i_pixel];
      if (d <= 0) {
        ray.ri = -1;
      } else {
        ray.r += d * ray.rd;
      }
    } else {
      ray.r += clip_near * ray.rd;
      d = scene_intersect(
        geometry_buffer,
        (const global uint *)render_buffer->texture_cdf,
        &ray,
        &texture_index);
    }
    if (i_bounce == 0) {
      ray.render_weight = (float3)(1.f, 1.f, 1.f);
      ray.pdf = 1.f;
    }
    if (ray.pdf < 1e-6f || 1.0f / ray.pdf < 1e-6f ||
        (ray.render_weight.x < 1e-6f && ray.render_weight.y < 1e-6f && ray.render_weight.z < 1e-6f)) {
      ray.ri = -2;
    } else {
      ray.render_weight /= ray.pdf;
      if (i_bounce > 1u)
        ray.render_last_weight /= ray.pdf;
    }
    const float3 envmap = material_envmap(render_buffer, envmap_index);
    if (ray.ri == -1) {
      ray.r += ray.rd;
      ray.rn = ray.rd;
      if (i_bounce >= 0u)
        ray.render += ray.render_weight * envmap;
    }
    material_normal(&ray, render_buffer, texture_index, i_bounce);
    if (ray.ri == -2) {
      ray.r += ray.rd;
      ray.pdf = 0.0f;
    }
    if (i_bounce <= 1u && i_sample < 128u) {
      intermediate_buffer->ray_r[0][i_bounce][i_pixel][i_sample] = ray.r.x;
      intermediate_buffer->ray_r[1][i_bounce][i_pixel][i_sample] = ray.r.y;
      intermediate_buffer->ray_r[2][i_bounce][i_pixel][i_sample] = ray.r.z;
      intermediate_buffer->ray_ri[i_bounce][i_pixel][i_sample] = ray.ri;
      intermediate_buffer->ray_pdf[i_bounce][i_pixel][i_sample] = ray.pdf;
    }
    if (ray.ri < 0)
      break;
    if (i_bounce == 0 && i_sample == 0) {
      render_buffer->ray[21][i_pixel] = ray.rt.x;
      render_buffer->ray[22][i_pixel] = ray.rt.y;
      render_buffer->ray[23][i_pixel] = clip_near + d;
      render_buffer->ray[24][i_pixel] = ray.rn.x;
      render_buffer->ray[25][i_pixel] = ray.rn.y;
      render_buffer->ray[26][i_pixel] = ray.rn.z;
    }
    if (i_bounce >= 0u) {
      const float3 intensity = material_intensity(render_buffer, texture_index, i_bounce);
      const float window = material_window(render_buffer, texture_index, i_bounce);
      ray.render += ray.render_weight * (intensity + window * envmap);
      if (i_bounce > 1u)
        ray.render_last += ray.render_last_weight * (intensity + window * envmap);
    }
    const float3 next_rd = sample_reflection(
      geometry_buffer,
      render_buffer,
      &ray,
      texture_index,
      i_bounce,
      low_discrepancy);
    const float3 brdf = diffuse_brdf(render_buffer, &ray, next_rd, texture_index, i_bounce) +
                        specular_brdf(render_buffer, &ray, next_rd, texture_index, i_bounce);
    ray.pdf = diffuse_pdf(&ray, next_rd) + specular_pdf(render_buffer, &ray, next_rd, texture_index, i_bounce);
    ray.rd = next_rd;
    ray.render_weight *= brdf;
    if (i_bounce >= 1u)
      ray.render_last_weight *= brdf;
  }
  if (i_sample < 128u) {
    intermediate_buffer->ray_render_last[0][i_pixel][i_sample] = ray.render_last.x;
    intermediate_buffer->ray_render_last[1][i_pixel][i_sample] = ray.render_last.y;
    intermediate_buffer->ray_render_last[2][i_pixel][i_sample] = ray.render_last.z;
  }
  local_gather_atomic_add_f(&render_buffer->ray[18][i_pixel], operand_gather, ray.render.x / 1024u);
  local_gather_atomic_add_f(&render_buffer->ray[19][i_pixel], operand_gather, ray.render.y / 1024u);
  local_gather_atomic_add_f(&render_buffer->ray[20][i_pixel], operand_gather, ray.render.z / 1024u);
}

kernel void kernel_ray_tracing_backward(
    const int srand,
    const uint low_discrepancy,
    const global GeometryBuffer *restrict geometry_buffer,
    volatile global RenderBuffer *restrict render_buffer,
    const global IntermediateBuffer *restrict intermediate_buffer) {
  local float operand_gather[128u];
  const uint i_pixel = get_global_id(0);
  const uint i_sample = get_global_id(1);
  float3 chain_render_weight[1u + 1];
  float3 chain_render[1u + 1];
  chain_render_weight[0] = (float3)(1.f, 1.f, 1.f);
  Ray ray = ray_init(srand, low_discrepancy, render_buffer);
  uint2 texture_index = (uint2)(0u, 0u);
  uint3 envmap_index = (uint3)(0u, 0u, 0u);
  for (int i_bounce = 0; i_bounce <= 1u; ++i_bounce) {
    if (ray.ri < 0) {
      chain_render_weight[i_bounce] = (float3)(0.f, 0.f, 0.f);
      chain_render[i_bounce] = (float3)(0.f, 0.f, 0.f);
      continue;
    }
    const float3 next_r = (float3)(intermediate_buffer->ray_r[0][i_bounce][i_pixel][i_sample],
                                   intermediate_buffer->ray_r[1][i_bounce][i_pixel][i_sample],
                                   intermediate_buffer->ray_r[2][i_bounce][i_pixel][i_sample]);
    float3 curr_rd = next_r - ray.r;
    curr_rd /= sqrt(dot(curr_rd, curr_rd));
    ray.ri = intermediate_buffer->ray_ri[i_bounce][i_pixel][i_sample];
    ray.pdf = intermediate_buffer->ray_pdf[i_bounce][i_pixel][i_sample];
    if (ray.ri == -2) {
      chain_render_weight[i_bounce] = (float3)(0.f, 0.f, 0.f);
      chain_render[i_bounce] = (float3)(0.f, 0.f, 0.f);
      continue;
    }
    if (i_bounce > 0) {
      const float3 brdf = diffuse_brdf(render_buffer, &ray, curr_rd, texture_index, i_bounce - 1) +
                          specular_brdf(render_buffer, &ray, curr_rd, texture_index, i_bounce - 1);
      chain_render_weight[i_bounce] = brdf / ray.pdf;
    }
    ray.rd = curr_rd;
    ray.r = next_r;
    background_intersect(render_buffer, &ray, &envmap_index);
    const float3 envmap = material_envmap(render_buffer, envmap_index);
    if (ray.ri == -1) {
    if (i_bounce >= 0u) {
      chain_render[i_bounce] = envmap;
    } else {
      chain_render[i_bounce] = (float3)(0.f, 0.f, 0.f);
    }
      continue;
    }
    if (ray.ri < 0) {
      chain_render_weight[i_bounce] = (float3)(0.f, 0.f, 0.f);
      chain_render[i_bounce] = (float3)(0.f, 0.f, 0.f);
      continue;
    }
    if (!(false && i_bounce == 0))
      update_ray(geometry_buffer, &ray, 0.f, ray.ri, &texture_index);
    material_normal(&ray, render_buffer, texture_index, i_bounce);
    if (i_bounce >= 0u) {
      const float3 intensity = material_intensity(render_buffer, texture_index, i_bounce);
      const float window = material_window(render_buffer, texture_index, i_bounce);
      chain_render[i_bounce] = intensity + window * envmap;
    } else {
      chain_render[i_bounce] = (float3)(0.f, 0.f, 0.f);
    }
  }
  chain_render[1u] += (float3)(intermediate_buffer->ray_render_last[0][i_pixel][i_sample],
                                                intermediate_buffer->ray_render_last[1][i_pixel][i_sample],
                                                intermediate_buffer->ray_render_last[2][i_pixel][i_sample]);
  for (int i_bounce = 1u - 1; i_bounce >= 0; --i_bounce)
    chain_render[i_bounce] += chain_render[i_bounce + 1] * chain_render_weight[i_bounce + 1];
  ray = ray_init(srand, low_discrepancy, render_buffer);
  float3 d_render = (float3)(render_buffer->ray[18][i_pixel],
                             render_buffer->ray[19][i_pixel],
                             render_buffer->ray[20][i_pixel]) / (float)128u;
  for (int i_bounce = 0; i_bounce <= 1u; ++i_bounce) {
    bool is_continue = false;
    if (ray.ri < 0) {
      d_render = (float3)(0.f, 0.f, 0.f);
      is_continue = true;
      goto d_continue_before_brdf;
    }
    const float3 next_r = (float3)(intermediate_buffer->ray_r[0][i_bounce][i_pixel][i_sample],
                                   intermediate_buffer->ray_r[1][i_bounce][i_pixel][i_sample],
                                   intermediate_buffer->ray_r[2][i_bounce][i_pixel][i_sample]);
    float3 curr_rd = next_r - ray.r;
    curr_rd /= sqrt(dot(curr_rd, curr_rd));
    ray.ri = intermediate_buffer->ray_ri[i_bounce][i_pixel][i_sample];
    ray.pdf = intermediate_buffer->ray_pdf[i_bounce][i_pixel][i_sample];
    float3 envmap = (float3)(0.f, 0.f, 0.f);
    if (ray.ri == -2) {
      d_render = (float3)(0.f, 0.f, 0.f);
      is_continue = true;
      goto d_continue_before_brdf;
    }
d_continue_before_brdf:
    if (i_bounce > 0) {
      const float3 d_brdf = is_continue ? (float3)(0.f, 0.f, 0.f) : d_render * chain_render[i_bounce] / ray.pdf;
      d_diffuse_brdf(render_buffer, operand_gather, &ray, curr_rd, texture_index, i_bounce - 1, d_brdf);
      d_specular_brdf(render_buffer, operand_gather, &ray, curr_rd, texture_index, i_bounce - 1, d_brdf);
      d_render *= chain_render_weight[i_bounce];
    }
    if (is_continue)
      goto d_continue_after_brdf;
    ray.rd = curr_rd;
    ray.r = next_r;
    background_intersect(render_buffer, &ray, &envmap_index);
    envmap = material_envmap(render_buffer, envmap_index);
    if (ray.ri == -1 && i_bounce >= 0u)
      d_material_envmap(render_buffer, envmap_index, d_render);
    if (ray.ri < 0) {
      d_render = (float3)(0.f, 0.f, 0.f);
      goto d_continue_after_brdf;
    }
    if (!(false && i_bounce == 0))
      update_ray(geometry_buffer, &ray, 0.f, ray.ri, &texture_index);
    material_normal(&ray, render_buffer, texture_index, i_bounce);
    const float window = material_window(render_buffer, texture_index, i_bounce);
    if (i_bounce >= 0u)
      d_material_envmap(render_buffer, envmap_index, window * d_render);
d_continue_after_brdf:
    if (i_bounce >= 0u) {
      d_material_intensity(render_buffer, operand_gather, texture_index, i_bounce, d_render);
      d_material_window(render_buffer, operand_gather, texture_index, i_bounce, dot(d_render, envmap));
    }
  }
}

void collect_texture(volatile global float (*data)[256u][256u][8u]) {
  volatile global float (*data_offset)[8u] = &(*data)[get_global_id(0)][get_global_id(1)];
  for (int i = 1; i < 8u; ++i)
    (*data_offset)[0] += (*data_offset)[i];
}

kernel void kernel_ray_tracing_backward_collect_texture(volatile global RenderBuffer *render_buffer) {
  for (int i = 0; i < 14; ++i)
    collect_texture(&render_buffer->d_texture[i]);
}

void collect_envmap(volatile global float (*data)[512u][512u][8u]) {
  volatile global float (*data_offset)[8u] = &(*data)[get_global_id(0)][get_global_id(1)];
  for (int i = 1; i < 8u; ++i)
    (*data_offset)[0] += (*data_offset)[i];
}

kernel void kernel_ray_tracing_backward_collect_envmap(volatile global RenderBuffer *render_buffer) {
  for (int i = 0; i < 3; ++i)
    for (int j = 0; j < 6; ++j)
      collect_envmap(&render_buffer->d_envmap[i][j]);
}

 

 

...全文
108 1 打赏 收藏 转发到动态 举报
AI 作业
写回复
用AI写文章
1 条回复
切换为时间正序
请发表友善的回复…
发表回复
  • 打赏
  • 举报
回复
谢谢反馈,已经请技术同事进行分析了

168

社区成员

发帖
与我相关
我的任务
社区描述
摩尔线程成立于 2020 年 10 月,以全功能 GPU 为核心,致力于向全球提供加速计算的基础设施和一站式解决方案,为各行各业的数智化转型提供强大的 AI 计算支持。 我们的目标是成为具备国际竞争力的 GPU 领军企业,为融合人工智能和数字孪生的数智世界打造先进的加速计算平台。我们的愿景是为美好世界加速。
企业社区
社区管理员
  • 摩尔线程
加入社区
  • 近7日
  • 近30日
  • 至今
社区公告
暂无公告

试试用AI创作助手写篇文章吧