未來研究方向

第六章結論與未來方向

6.2 未來研究方向

由於本次的設計主要還是在於 OpenCL 記憶體使用限制下的場景，未來應可以設計針對當資料量極大時所採取的切割以及分流方法，但是依照目前為止 CPU 廠商則盡可能希望的 OpenCL 發展偏向直接使用主記憶體或是與 C 語言共用資料結構，也可能得以解決這些問題。

而光追蹤演算法內也可透過變更演算法的設計以致於其他適合平行化的部分的方法，甚至可以加入次要光線追蹤的效果，這些是未來可能的研究方向。

附錄

1. Intersect OpenCL 程式碼

void triangle4_intersect(__global struct Triangle4 *triangle4, struct Ray *ray, __global struct Hit *hit)

{

float3 O = (float3) (ray->org[0], ray->org[1], ray->org[2]);

float3 D = (float3) (ray->dir[0], ray->dir[1], ray->dir[2]);

float4 C[3] = {triangle4->v0[0]-O.x, triangle4->v0[1]-O.y, triangle4->v0[2]-O.z};

float4 R[3] = SSE3F_FLOAT3_CROSS(C, -D);

float4 det = SSE3F_FLOAT3_DOT(triangle4->Ng, D);

float4 T = SSE3F_SSE3F_DOT(triangle4->Ng, C);

float4 U = SSE3F_SSE3F_DOT(R, triangle4->e2);

float4 V = SSE3F_SSE3F_DOT(R, triangle4->e1);

float4 absDet = ABS(det);

float4 signDet = select((float4)1.0, (float4)-1.0, signbit(det));

float4 _t = T * signDet; float4 _u = U * signDet; float4 _v = V * signDet;

float4 _w = absDet - _u - _v;

int4 mask = (triangle4->id0 != -1) & (det != 0.f) & (_t >= absDet*ray->near) & (absDet*hit->t >=

_t) & (min(min(_u,_v),_w) >= 0.f);

if (! any(mask)) return;

float4 rcpAbsDet = 1.0f / absDet;

float4 t = _t * rcpAbsDet; float4 u = _u * rcpAbsDet; float4 v = _v * rcpAbsDet;

float4 __t = select(INFINITY, t, mask);

float4 tmp = min(__t, __t.zwxy);

float4 reducedMin = min(tmp, tmp.yxwz);

mask = isequal(__t, reducedMin);

int tri = round(dot(select((float4)0, (float4)(0,1,2,3), mask), (float4)1.0));

float tA[4] = {t.x, t.y, t.z, t.w};

float uA[4] = {u.x, u.y, u.z, u.w};

float vA[4] = {v.x, v.y, v.z, v.w};

int id0A[4] = {triangle4->id0.x, triangle4->id0.y, triangle4->id0.z, triangle4->id0.w};

int id1A[4] = {triangle4->id1.x, triangle4->id1.y, triangle4->id1.z, triangle4->id1.w};

hit->t = tA[tri]; hit->u = uA[tri]; hit->v = vA[tri];

hit->id0 = id0A[tri];

hit->id1 = id1A[tri];

return;

}

__kernel void BVH4Traverse(__global struct Ray *rays, __global struct Hit *hits,__global int

*bvh_root, __global struct BVH4_Node *bvh_nodes, __global struct Triangle4 *bvh_triangles) {

__global struct BVH4_Node *p = bvh_nodes;

int idx = get_global_id(0); int nodeID, ofs, num; struct Ray ray = rays[idx];

float4 tNearX, tNearY, tNearZ, tNear, tFarX, tFarY, tFarZ, tFar;

int4 _hit; int numHits=0; int hitList[4];

int stackPtr = 1; int stack[100]; stack[0] = *bvh_root;

const int nearX = ray.dir[0] >= 0 ? 0*sizeof(float4) : 1*sizeof(float4);

const int nearY = ray.dir[1] >= 0 ? 2*sizeof(float4) : 3*sizeof(float4);

const int nearZ = ray.dir[2] >= 0 ? 4*sizeof(float4) : 5*sizeof(float4);

const int farX = nearX ^ sizeof(float4); const int farY = nearY ^ sizeof(float4);

const int farZ = nearZ ^ sizeof(float4); hits[idx].t = ray.far;

while (stackPtr > 0) {

nodeID = stack[--stackPtr];

if (nodeID >= 0) {

p = node(bvh_nodes, nodeID);

tNearX = (float4_ITEM(p,nearX) - ray.org[0]) * ray.rdir[0];

tNearY = (float4_ITEM(p,nearY) - ray.org[1]) * ray.rdir[1];

tNearZ = (float4_ITEM(p,nearZ) - ray.org[2]) * ray.rdir[2];

tNear = max( max( max(tNearX, tNearY), tNearZ ), ray.near );

tFarX = (float4_ITEM(p,farX) - ray.org[0]) * ray.rdir[0];

tFarY = (float4_ITEM(p,farY) - ray.org[1]) * ray.rdir[1];

tFarZ = (float4_ITEM(p,farZ) - ray.org[2]) * ray.rdir[2];

tFar = min( min( min(tFarX, tFarY), tFarZ ), ray.far );

_hit = tNear <= tFar;

int _hitA[4] = {_hit.x, _hit.y, _hit.z, _hit.w};

float tNearA[4] = {tNear.x, tNear.y, tNear.z, tNear.w};

numHits = 0;

for (int i=0; i<4; i++) {

if (_hitA[i]) hitList[ numHits++ ] = i;

}

if (0 == numHits) continue;

if (1 == numHits) {

stack[stackPtr++] = p->child[ hitList[0] ];

continue;

}

if (2 == numHits) {

if (tNearA[hitList[0]] < tNearA[hitList[1]]) {

stack[stackPtr++] = p->child[ hitList[1] ];

stack[stackPtr++] = p->child[ hitList[0] ];

} else {

stack[stackPtr++] = p->child[ hitList[0] ];

stack[stackPtr++] = p->child[ hitList[1] ];

}

int dMax = tNearA[hitList[i]];

for (int j=i+1; j<numHits; j++) {

hitList[cMax] = tmp;

}

stack[stackPtr++] = p->child[ hitList[i] ];

} // for i } // if (nodeID >= 0) else {

nodeID ^= 0x80000000; ofs = nodeID >> 5; num = nodeID & 0x1F;

if (!num) continue;

for (int i=ofs; i<ofs+num; i++) {

triangle4_intersect( &bvh_triangles[i], &ray, &hits[idx] );

}

ray.far = hits[idx].t;

} // else

} // while (stackPtr > 0) }

2. Occluded OpenCL 程式碼

int triangle4_occluded(__global struct Triangle4 *triangle4, struct Ray *ray)

{

float3 O = (float3) (ray->org[0], ray->org[1], ray->org[2]);

float3 D = (float3) (ray->dir[0], ray->dir[1], ray->dir[2]);

float4 C[3] = {triangle4->v0[0]-O.x, triangle4->v0[1]-O.y, triangle4->v0[2]-O.z};

float4 R[3] = SSE3F_FLOAT3_CROSS(C, -D);

float4 det = SSE3F_FLOAT3_DOT(triangle4->Ng, D);

float4 T = SSE3F_SSE3F_DOT(triangle4->Ng, C);

float4 U = SSE3F_SSE3F_DOT(R, triangle4->e2);

float4 V = SSE3F_SSE3F_DOT(R, triangle4->e1);

float4 absDet = ABS(det);

float4 signDet = select((float4)1.0, (float4)-1.0, signbit(det));

float4 _t = T * signDet;

float4 _u = U * signDet;

float4 _v = V * signDet;

float4 _w = absDet - _u - _v;

int4 mask = (triangle4->id0 != -1) & (det != 0.f) & (_t >= absDet*ray->near) & (absDet*ray->far

>= _t) & (min(min(_u,_v),_w) >= 0.f);

return any(mask);

}

__kernel void BVH4Occluded(__global struct Ray *rays, __global int *ishit, __global int

*bvh_root, __global struct BVH4_Node *bvh_nodes, __global struct Triangle4 *bvh_triangles) {

__global struct BVH4_Node *p = bvh_nodes;

int idx = get_global_id(0);

ishit[idx]=0;

int nodeID, ofs, num;

struct Ray ray = rays[idx];

float4 tNearX, tNearY, tNearZ, tNear, tFarX, tFarY, tFarZ, tFar;

int4 _hit;

const int farX = nearX ^ sizeof(float4); const int farY = nearY ^ sizeof(float4);

const int farZ = nearZ ^ sizeof(float4);

while (stackPtr > 0) {

nodeID = stack[--stackPtr];

if (nodeID >= 0) {

p = node(bvh_nodes, nodeID);

tNearX = (float4_ITEM(p,nearX) - ray.org[0]) * ray.rdir[0];

tNearY = (float4_ITEM(p,nearY) - ray.org[1]) * ray.rdir[1];

tNearZ = (float4_ITEM(p,nearZ) - ray.org[2]) * ray.rdir[2];

tNear = max( max( max(tNearX, tNearY), tNearZ ), ray.near );

tFarX = (float4_ITEM(p,farX) - ray.org[0]) * ray.rdir[0];

tFarY = (float4_ITEM(p,farY) - ray.org[1]) * ray.rdir[1];

tFarZ = (float4_ITEM(p,farZ) - ray.org[2]) * ray.rdir[2];

tFar = min( min( min(tFarX, tFarY), tFarZ ), ray.far );

_hit = tNear <= tFar;

int _hitA[4] = {_hit.x, _hit.y, _hit.z, _hit.w};

numHits = 0;

for (int i=0; i<4; i++) { if (_hitA[i]) {

stack[stackPtr++] = p->child[i];

} }

} // if (nodeID >= 0) else {

nodeID ^= 0x80000000;

ofs = nodeID >> 5;

num = nodeID & 0x1F;

if (!num) continue;

for (int i=ofs; i<ofs+num; i++) {

if(triangle4_occluded(&bvh_triangles[i], &ray)){

ishit[idx]=1; return;

}

參考文獻

參考書目：

[Du12] P. Du, R. Weber, P. Luszczek, S. Tomov, G. Peterson, J. Dongarra:

From CUDA to OpenCL: Towards a Performance-portable Solution for Multi-platform GPU Programming. Elsevier Science Publishers B. V. Amsterdam, 2012.

[GHKMS11] B. Gaster L. Howes, D. R. Kaeli, P. Mistry, D. Schaa:

Heterogeneous Computing with OpenCL. Morgan Kaufmann, 2011.

[KH12] R. Karrenberg and S. Hack:Improving Performance of OpenCL on CPUs. Springer-Verlag Berlin, 2012.

[MGMF11] A. Munshi, B. Gaster, T. G. Mattson, J. Fung, D. Ginsburg:

OpenCL Programming Guide. Addison-Wesley, 2011.

[Rei08] S. Reiter: Real–time Ray Tracing of Dynamic Scenes. Diploma Thesis, Institute for Graphics and Parallel Processing, Johannes Kepler University, 2008.

[SK11] J. Sanders, E. Kandrot: Cuda by Example:An Introduction to General-purpose GPU Programming. Addison-Wesley, 2011.

[Sca11] M. Scarpino: OpenCL in Action: How to Accelerate Graphics and Computations. Manning Publications, 2011.

[Shi09] M. Shih, Y.-F. Chiu, Y.-C. Chen, C.-F. Chang: Real-Time Ray Tracing with CUDA. In ICA3PP ’09: Proceedings of the 9th International Conference on Algorithms and Architectures for Parallel

Processing, pp. 327–337, 2009.

[Suf07] K. Suffern: Ray Tracing from the Ground Up. A K Peters , 2007.

[Wat99] A. Watt: 3D Computer Graphics.Addison Wesley, 1999.

[Woo04] S. Woop: A Ray Tracing Hardware Architecture for Dynamic Scenes. Computer Graphics Lab, Saarland University, 2004.

參考網站：

[1.] 經典光跡追蹤演算法

http://zh.wikipedia.org/wiki/%E5%85%89%E7%B7%9A%E8%BF%BD%E8%B9%

A4

[2.]OpenCL Introduction

http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/

[3.]Intel Ivy Bridge

http://en.wikipedia.org/wiki/Ivy_Bridge_%28microarchitecture%29

[4.]NVidia CUDA

http://en.wikipedia.org/wiki/CUDA

[5.]Intel OpenCL Optimization-Guide

http://software.intel.com/sites/landingpage/opencl/optimization-guide/

[6.] Intel Streaming_SIMD_Extensions

http://en.wikipedia.org/wiki/Streaming_SIMD_Extensions

[7.] OpenCL Programming for the CUDA Architecture

http://www.nvidia.com/content/cudazone/download/OpenCL/NVIDIA_OpenCL_Prog rammingOverview.pdf

[8.] CUDA Pinned memory

http://www.orangeowlsolutions.com/archives/443

在文檔中優化在異構性平台的OpenCL程式撰寫架構 -以光追蹤為例 (頁 48-55)

第六章 結論與未來方向

6.2 未來研究方向

附 錄

參考文獻

參考書目：

參考網站：

http://zh.wikipedia.org/wiki/%E5%85%89%E7%B7%9A%E8%BF%BD%E8%B9%

A4

http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/

http://en.wikipedia.org/wiki/Ivy_Bridge_%28microarchitecture%29

http://en.wikipedia.org/wiki/CUDA

http://software.intel.com/sites/landingpage/opencl/optimization-guide/

http://en.wikipedia.org/wiki/Streaming_SIMD_Extensions

http://www.nvidia.com/content/cudazone/download/OpenCL/NVIDIA_OpenCL_Prog rammingOverview.pdf

http://www.orangeowlsolutions.com/archives/443

第六章結論與未來方向

附錄