3 Replies Latest reply on Dec 18, 2009 8:21 AM by genaganna

    Compilation error : "declared implicitly"

    spectral

      Hi,

      In my OpenCL code I need to call a method that will do a specific processing on a "set of datas".

      So I have create a method like this :

      int Intersects_Primitive(float * datas, IntersectionResult * result)

      {

      ...

      }

       

      But when I try to call it, I get a compilation error !! I'm sure that there is a reason, but which one ? and what solution to this problem ?

      Intersects_Primitive(&datas, &result)

      I got the following message "declared implicitly"... what does it mean ?

      Thanks

        • Compilation error : "declared implicitly"
          genaganna

          could you please post kernel code to reproduce this error so that we can look for workaround?

            • Compilation error : "declared implicitly"
              spectral

              Hi,

              Here is my code :

               

               

              typedef struct

              {

              /// <summary>

              /// The bounding box's coordinates

              /// </summary>

              float MinX, MinY, MinZ, MaxX, MaxY, MaxZ;

              } AABB;

              typedef struct

              {

              float Maximum;

              int PrimitiveId;

              float U;

              float V;

              } IntersectionResult;

              typedef struct

              {

              float x, y, z, w;

              } Vector3;

              typedef struct

              {

              /// <summary>

              /// The bounding box's coordinates

              /// </summary>

              AABB BBox;

              /// <summary>

              /// The index of the sibbling node (next node at the same level).

              /// </summary>

              int SkipNodeIndex;

              /// <summary>

              /// The primitive Id.

              /// </summary>

              /// <remarks>If it is a set (Count > -1) it is the first index in the primitives' ids list (_primitivesIds)</remarks>

              int PrimitiveId;

              /// <summary>

              /// The number of primitives, used when it is a set.

              /// </summary>

              ushort Count;

              } BoundingVolume;

              kernel void TreeTraversal(

              global Vector3 * origin,

              global Vector3 * direction,

              float rayMaximum,

              int isShadowRay, // 1 = true, 0 = false

              global Vector3 * rayOffset,

              global BoundingVolume * _nodes,

              global int * primitivesIds,

              global Vector3 * primitivesVertices,

              global IntersectionResult * result)

              {

              int OCLIndex = get_global_id(0);

               

              int bvNodeIndex = 0;

              //MITData mitData = new MITData(ray);

              // End of the tree

              int stopNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex;

              int hasIntersection = 0;

               

              int loopCount = 0;

              while (bvNodeIndex < stopNodeIndex)

              {

              loopCount++;

              // Test for the best primitives

              //if (loopCount < 1 /*&& hitsCache != null*/)

              //{

              // int[] bestHits = hitsCache.Primitives;

              // if (bestHits != null)

              // for (int index = 0; index < bestHits.Length; index++)

              // _primitiveList.IntersectPrimitive(ray, bestHits[index], istate);

              //}

               

              // Do a hit test with the bounding volume.

              // We use the "closest intersection" to check against the BV intersection.

              // If the 'closest intersection' < 'bv intersection' then there is no

              // primitive intersection possible !

              float minHit;

              float maxHit;

              int hasHit = Intersects_BoxRay(/*mitData,*/ origin, direction, _nodes[bvNodeIndex].BBox, &minHit, &maxHit);

              if (hasHit &&

              minHit <= rayMaximum)

              {

              // It is a leaf -> test the primitives

              if (_nodes[bvNodeIndex].PrimitiveId > -1)

              {

              origin = origin + rayOffset;

              // Contains a simple primitive

              if (_nodes[bvNodeIndex].Count < 1)

              {

              if (Intersects_Primitive(primitivesVertices, origin, direction, _nodes[bvNodeIndex].PrimitiveId, result))

              {

              // If shadow ray

              if (isShadowRay > 0)

              return true;

              //if (hitsCache != null)

              // hitsCache.AddHit(_nodes[bvNodeIndex].PrimitiveId, ray.Maximum);

              hasIntersection = 1;

              }

              }

              // Contains a set of primitive

              else if (_nodes[bvNodeIndex].Count > 0)

              {

              int startIndex = bv.PrimitiveId;

              int endIndex = startIndex + bv.Count - 1;

              for (int index = startIndex; index <= endIndex; index++)

              if (Intersects_Primitive(primitivesVertices, origin, direction, _primitivesIds[index], result))

              {

              // If shadow ray

              if (isShadowRay > 0)

              return true;

              //if (hitsCache != null)

              // hitsCache.AddHit(_primitivesIds[index], ray.Maximum);

              hasIntersection = 1;

              }

              }

              origin -= rayOffset;

              }

              // Next node at the same level OR

              // the next sibbling of the parent.

              bvNodeIndex++;

              if ((bvNodeIndex >= stopNodeIndex || bvNodeIndex == _nodes[bvNodeIndex].SkipNodeIndex) &&

              hasIntersection)

              return 1;

              }

              // Continue at the same level

              else

              bvNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex;

              }

              return 0;

              }

              int Intersects_BoxRay(struct Vector3 * origin, struct Vector3 * direction, struct BoundingVolume * aabb, float * minHit, float * maxHit)

              {

              return 0;

              }

              int Intersects_Primitive(float * primitivesVertices, struct Vector3 * origin, struct Vector3 * direction, int * _primitivesIds, struct IntersectionResult * result)

              {

              return 0;

              }



                • Compilation error : "declared implicitly"
                  genaganna

                  OpenCL is based on c99 with some extensions and restrictions.

                   

                  1. functions should be defined before invoking

                  2. pointer casting has lot of restrictions. see 6.8 section in OpenCL spec

                  3. user is responsible for arithmatic for user defined variables.

                  4. There are restrictions on structure alignments. see release notes

                   

                  I removed all errors from your code. see attached code.

                    Note : Initialized dummy values for undefined variables.

                  typedef struct { /// <summary> /// The bounding box's coordinates /// </summary> float MinX, MinY, MinZ, MaxX, MaxY, MaxZ; } AABB; typedef struct { float Maximum; int PrimitiveId; float U; float V; }IntersectionResult; typedef struct { float x, y, z, w; } Vector3; typedef struct { /// <summary> /// The bounding box's coordinates /// </summary> AABB BBox; /// <summary> /// The index of the sibbling node (next node at the same level). /// </summary> int SkipNodeIndex; /// <summary> /// The primitive Id. /// </summary> /// <remarks>If it is a set (Count > -1) it is the first index in the primitives' ids list (_primitivesIds)</remarks> int PrimitiveId; /// <summary> /// The number of primitives, used when it is a set. /// </summary> ushort Count; }BoundingVolume; int Intersects_BoxRay(global Vector3 * origin, global Vector3 * direction, global AABB* aabb, float * minHit, float * maxHit) { return 0; } int Intersects_Primitive(global Vector3 * primitivesVertices, global Vector3 * origin, global Vector3 * direction, int _primitivesIds, global IntersectionResult * result) { return 0; } kernel void TreeTraversal( global Vector3 * origin, global Vector3 * direction, float rayMaximum, int isShadowRay, // 1 = true, 0 = false global Vector3 * rayOffset, global BoundingVolume * _nodes, global int * primitivesIds, global Vector3 * primitivesVertices, global IntersectionResult * result) { int OCLIndex = get_global_id(0); int bvNodeIndex = 0; //MITData mitData = new MITData(ray); // End of the tree int stopNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex; int hasIntersection = 0; int loopCount = 0; while (bvNodeIndex < stopNodeIndex) { loopCount++; // Test for the best primitives //if (loopCount < 1 /*&& hitsCache != null*/) //{ // int[] bestHits = hitsCache.Primitives; // if (bestHits != null) // for (int index = 0; index < bestHits.Length; index++) // _primitiveList.IntersectPrimitive(ray, bestHits[index], istate); //} // Do a hit test with the bounding volume. // We use the "closest intersection" to check against the BV intersection. // If the 'closest intersection' < 'bv intersection' then there is no // primitive intersection possible ! float minHit; float maxHit; int hasHit = Intersects_BoxRay(/*mitData,*/ origin, direction, &_nodes[bvNodeIndex].BBox, &minHit, &maxHit); if (hasHit && minHit <= rayMaximum) { // It is a leaf -> test the primitives if (_nodes[bvNodeIndex].PrimitiveId > -1) { origin[0].x = origin[0].x + rayOffset[0].x; origin[0].y = origin[0].y + rayOffset[0].y; origin[0].z = origin[0].z + rayOffset[0].z; //origin = origin + rayOffset; // Contains a simple primitive if(_nodes[bvNodeIndex].Count < 1) { if(Intersects_Primitive(primitivesVertices, origin, direction, _nodes[bvNodeIndex].PrimitiveId, result)) { // If shadow ray if (isShadowRay > 0) return; //if (hitsCache != null) // hitsCache.AddHit(_nodes[bvNodeIndex].PrimitiveId, ray.Maximum); hasIntersection = 1; } } // Contains a set of primitive else if (_nodes[bvNodeIndex].Count > 0) { //Note : bv variable not defined int startIndex = 0;//bv.PrimitiveId; int endIndex = startIndex + 0;//bv.Count - 1; for (int index = startIndex; index <= endIndex; index++) if(Intersects_Primitive(primitivesVertices, origin, direction, primitivesIds[index], result)) { // If shadow ray if (isShadowRay > 0) return; //if (hitsCache != null) // hitsCache.AddHit(_primitivesIds[index], ray.Maximum); hasIntersection = 1; } } //Note : Arithmetic on user defined variables //origin[0] -= rayOffset origin[0].x -= rayOffset[0].x; origin[0].y -= rayOffset[0].y; origin[0].z -= rayOffset[0].z; } // Next node at the same level OR // the next sibbling of the parent. bvNodeIndex++; if ((bvNodeIndex >= stopNodeIndex || bvNodeIndex == _nodes[bvNodeIndex].SkipNodeIndex) && hasIntersection) return; } // Continue at the same level else bvNodeIndex = _nodes[bvNodeIndex].SkipNodeIndex; } return; }