nested classes with pointers woth openACC












0














I have a rather big code in C++, I had to integrate some new class to the base class as shown below.



class A
{
int N;
B b;
double *__restrict__ w;
construct();
}



A::construct()
{
w=new double[N];

#pragma acc data enter create(this)
#pragma acc update device(this)
#pragma acc data enter create(w)
// allocate class A
b.construct()
}



class B
{
double *__restrict__ u;
double *__restrict__ v;
B(){};
construct();
}
B::construct()
{
u=new double[N];
v=new double[N];
#pragma acc data enter create(this)
#pragma acc update device(this)
#pragma acc data enter create(u)
#pragma acc data enter create(v)

}


I think I am running into the deep copy issue as the pointers of class B are invalidated and hence the behavior of the code on GPU i undefined.
I would appreciate the feedback on how to perform the class inclusion in another class without getting into the deep copy issue. I suspect the update device (this) somehow causes this.










share|improve this question



























    0














    I have a rather big code in C++, I had to integrate some new class to the base class as shown below.



    class A
    {
    int N;
    B b;
    double *__restrict__ w;
    construct();
    }



    A::construct()
    {
    w=new double[N];

    #pragma acc data enter create(this)
    #pragma acc update device(this)
    #pragma acc data enter create(w)
    // allocate class A
    b.construct()
    }



    class B
    {
    double *__restrict__ u;
    double *__restrict__ v;
    B(){};
    construct();
    }
    B::construct()
    {
    u=new double[N];
    v=new double[N];
    #pragma acc data enter create(this)
    #pragma acc update device(this)
    #pragma acc data enter create(u)
    #pragma acc data enter create(v)

    }


    I think I am running into the deep copy issue as the pointers of class B are invalidated and hence the behavior of the code on GPU i undefined.
    I would appreciate the feedback on how to perform the class inclusion in another class without getting into the deep copy issue. I suspect the update device (this) somehow causes this.










    share|improve this question

























      0












      0








      0







      I have a rather big code in C++, I had to integrate some new class to the base class as shown below.



      class A
      {
      int N;
      B b;
      double *__restrict__ w;
      construct();
      }



      A::construct()
      {
      w=new double[N];

      #pragma acc data enter create(this)
      #pragma acc update device(this)
      #pragma acc data enter create(w)
      // allocate class A
      b.construct()
      }



      class B
      {
      double *__restrict__ u;
      double *__restrict__ v;
      B(){};
      construct();
      }
      B::construct()
      {
      u=new double[N];
      v=new double[N];
      #pragma acc data enter create(this)
      #pragma acc update device(this)
      #pragma acc data enter create(u)
      #pragma acc data enter create(v)

      }


      I think I am running into the deep copy issue as the pointers of class B are invalidated and hence the behavior of the code on GPU i undefined.
      I would appreciate the feedback on how to perform the class inclusion in another class without getting into the deep copy issue. I suspect the update device (this) somehow causes this.










      share|improve this question













      I have a rather big code in C++, I had to integrate some new class to the base class as shown below.



      class A
      {
      int N;
      B b;
      double *__restrict__ w;
      construct();
      }



      A::construct()
      {
      w=new double[N];

      #pragma acc data enter create(this)
      #pragma acc update device(this)
      #pragma acc data enter create(w)
      // allocate class A
      b.construct()
      }



      class B
      {
      double *__restrict__ u;
      double *__restrict__ v;
      B(){};
      construct();
      }
      B::construct()
      {
      u=new double[N];
      v=new double[N];
      #pragma acc data enter create(this)
      #pragma acc update device(this)
      #pragma acc data enter create(u)
      #pragma acc data enter create(v)

      }


      I think I am running into the deep copy issue as the pointers of class B are invalidated and hence the behavior of the code on GPU i undefined.
      I would appreciate the feedback on how to perform the class inclusion in another class without getting into the deep copy issue. I suspect the update device (this) somehow causes this.







      gpu openacc pgi pgi-accelerator






      share|improve this question













      share|improve this question











      share|improve this question




      share|improve this question










      asked Dec 27 '18 at 16:23









      BatiCode

      189113




      189113
























          1 Answer
          1






          active

          oldest

          votes


















          0














          Do you have a full example which recreates the error you're seeing? I wrote a little test example using your code snip-it and it worked fine. (See below)



          If you were updating the "this" pointer after creating the arrays, then it would be a problem since you'd be overwriting the device pointers with the host pointers. But as you show above, it shouldn't be an issue.



          % cat test.cpp
          #include <iostream>

          class B
          {
          public:
          int N;
          double *__restrict__ u;
          double *__restrict__ v;
          void construct(int);
          };

          void B::construct(int _N)
          {
          N=_N;
          u=new double[N];
          v=new double[N];
          #pragma acc enter data create(this)
          #pragma acc update device(this)
          #pragma acc enter data create(u[:N])
          #pragma acc enter data create(v[:N])
          }

          class A
          {
          public:
          int N;
          B b;
          double *__restrict__ w;
          void construct(int);
          };

          void A::construct(int _N)
          {
          N=_N;
          w=new double[N];
          #pragma acc enter data create(this)
          #pragma acc update device(this)
          #pragma acc enter data create(w[:N])

          // allocate class A
          b.construct(N);
          }


          int main() {

          A myA;
          int N=32;
          myA.construct(N);

          #pragma acc parallel loop present(myA)
          for (int i=0; i<N; ++i) {
          myA.w[i] = i;
          myA.b.u[i] = i;
          myA.b.v[i] = i;
          }
          #pragma acc update host( myA.w[:N], myA.b.u[:N], myA.b.v[:N])
          for (int i=0; i<N; ++i) {
          std::cout << myA.w[i] << ":" << myA.b.u[i] << ":" << myA.b.v[i] << std::endl;
          }
          return 0;
          }
          % pgc++ test.cpp -Minfo=accel -V18.10 -ta=tesla; a.out
          main:
          49, Generating present(myA)
          Accelerator kernel generated
          Generating Tesla code
          52, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
          56, Generating update self(myA.b.u[:N],myA.w[:N],myA.b.v[:N])
          B::construct(int):
          21, Generating update device(this[:1])
          Generating enter data create(this[:1],v[:N],u[:N])
          A::construct(int):
          41, Generating update device(this[:1])
          Generating enter data create(w[:N],this[:1])
          0:0:0
          1:1:1
          2:2:2
          3:3:3
          4:4:4
          5:5:5
          6:6:6
          7:7:7
          8:8:8
          9:9:9
          10:10:10
          11:11:11
          12:12:12
          13:13:13
          14:14:14
          15:15:15
          16:16:16
          17:17:17
          18:18:18
          19:19:19
          20:20:20
          21:21:21
          22:22:22
          23:23:23
          24:24:24
          25:25:25
          26:26:26
          27:27:27
          28:28:28
          29:29:29
          30:30:30
          31:31:31





          share|improve this answer





















          • yes, Can I send the code to you ? Thanks
            – BatiCode
            Dec 27 '18 at 23:41










          • Sure. Send it to PGI Customer Service (trs@pgroup,com) and ask Alex to forward it to me.
            – Mat Colgrove
            Dec 28 '18 at 15:25










          • Thank you so much as I have been struggling with why this situation arises, as both classes are developed separately and they work fine. when I put them together the issue arises. I just submitted it. Thanks again for your time and help
            – BatiCode
            Dec 28 '18 at 16:55











          Your Answer






          StackExchange.ifUsing("editor", function () {
          StackExchange.using("externalEditor", function () {
          StackExchange.using("snippets", function () {
          StackExchange.snippets.init();
          });
          });
          }, "code-snippets");

          StackExchange.ready(function() {
          var channelOptions = {
          tags: "".split(" "),
          id: "1"
          };
          initTagRenderer("".split(" "), "".split(" "), channelOptions);

          StackExchange.using("externalEditor", function() {
          // Have to fire editor after snippets, if snippets enabled
          if (StackExchange.settings.snippets.snippetsEnabled) {
          StackExchange.using("snippets", function() {
          createEditor();
          });
          }
          else {
          createEditor();
          }
          });

          function createEditor() {
          StackExchange.prepareEditor({
          heartbeatType: 'answer',
          autoActivateHeartbeat: false,
          convertImagesToLinks: true,
          noModals: true,
          showLowRepImageUploadWarning: true,
          reputationToPostImages: 10,
          bindNavPrevention: true,
          postfix: "",
          imageUploader: {
          brandingHtml: "Powered by u003ca class="icon-imgur-white" href="https://imgur.com/"u003eu003c/au003e",
          contentPolicyHtml: "User contributions licensed under u003ca href="https://creativecommons.org/licenses/by-sa/3.0/"u003ecc by-sa 3.0 with attribution requiredu003c/au003e u003ca href="https://stackoverflow.com/legal/content-policy"u003e(content policy)u003c/au003e",
          allowUrls: true
          },
          onDemand: true,
          discardSelector: ".discard-answer"
          ,immediatelyShowMarkdownHelp:true
          });


          }
          });














          draft saved

          draft discarded


















          StackExchange.ready(
          function () {
          StackExchange.openid.initPostLogin('.new-post-login', 'https%3a%2f%2fstackoverflow.com%2fquestions%2f53947984%2fnested-classes-with-pointers-woth-openacc%23new-answer', 'question_page');
          }
          );

          Post as a guest















          Required, but never shown

























          1 Answer
          1






          active

          oldest

          votes








          1 Answer
          1






          active

          oldest

          votes









          active

          oldest

          votes






          active

          oldest

          votes









          0














          Do you have a full example which recreates the error you're seeing? I wrote a little test example using your code snip-it and it worked fine. (See below)



          If you were updating the "this" pointer after creating the arrays, then it would be a problem since you'd be overwriting the device pointers with the host pointers. But as you show above, it shouldn't be an issue.



          % cat test.cpp
          #include <iostream>

          class B
          {
          public:
          int N;
          double *__restrict__ u;
          double *__restrict__ v;
          void construct(int);
          };

          void B::construct(int _N)
          {
          N=_N;
          u=new double[N];
          v=new double[N];
          #pragma acc enter data create(this)
          #pragma acc update device(this)
          #pragma acc enter data create(u[:N])
          #pragma acc enter data create(v[:N])
          }

          class A
          {
          public:
          int N;
          B b;
          double *__restrict__ w;
          void construct(int);
          };

          void A::construct(int _N)
          {
          N=_N;
          w=new double[N];
          #pragma acc enter data create(this)
          #pragma acc update device(this)
          #pragma acc enter data create(w[:N])

          // allocate class A
          b.construct(N);
          }


          int main() {

          A myA;
          int N=32;
          myA.construct(N);

          #pragma acc parallel loop present(myA)
          for (int i=0; i<N; ++i) {
          myA.w[i] = i;
          myA.b.u[i] = i;
          myA.b.v[i] = i;
          }
          #pragma acc update host( myA.w[:N], myA.b.u[:N], myA.b.v[:N])
          for (int i=0; i<N; ++i) {
          std::cout << myA.w[i] << ":" << myA.b.u[i] << ":" << myA.b.v[i] << std::endl;
          }
          return 0;
          }
          % pgc++ test.cpp -Minfo=accel -V18.10 -ta=tesla; a.out
          main:
          49, Generating present(myA)
          Accelerator kernel generated
          Generating Tesla code
          52, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
          56, Generating update self(myA.b.u[:N],myA.w[:N],myA.b.v[:N])
          B::construct(int):
          21, Generating update device(this[:1])
          Generating enter data create(this[:1],v[:N],u[:N])
          A::construct(int):
          41, Generating update device(this[:1])
          Generating enter data create(w[:N],this[:1])
          0:0:0
          1:1:1
          2:2:2
          3:3:3
          4:4:4
          5:5:5
          6:6:6
          7:7:7
          8:8:8
          9:9:9
          10:10:10
          11:11:11
          12:12:12
          13:13:13
          14:14:14
          15:15:15
          16:16:16
          17:17:17
          18:18:18
          19:19:19
          20:20:20
          21:21:21
          22:22:22
          23:23:23
          24:24:24
          25:25:25
          26:26:26
          27:27:27
          28:28:28
          29:29:29
          30:30:30
          31:31:31





          share|improve this answer





















          • yes, Can I send the code to you ? Thanks
            – BatiCode
            Dec 27 '18 at 23:41










          • Sure. Send it to PGI Customer Service (trs@pgroup,com) and ask Alex to forward it to me.
            – Mat Colgrove
            Dec 28 '18 at 15:25










          • Thank you so much as I have been struggling with why this situation arises, as both classes are developed separately and they work fine. when I put them together the issue arises. I just submitted it. Thanks again for your time and help
            – BatiCode
            Dec 28 '18 at 16:55
















          0














          Do you have a full example which recreates the error you're seeing? I wrote a little test example using your code snip-it and it worked fine. (See below)



          If you were updating the "this" pointer after creating the arrays, then it would be a problem since you'd be overwriting the device pointers with the host pointers. But as you show above, it shouldn't be an issue.



          % cat test.cpp
          #include <iostream>

          class B
          {
          public:
          int N;
          double *__restrict__ u;
          double *__restrict__ v;
          void construct(int);
          };

          void B::construct(int _N)
          {
          N=_N;
          u=new double[N];
          v=new double[N];
          #pragma acc enter data create(this)
          #pragma acc update device(this)
          #pragma acc enter data create(u[:N])
          #pragma acc enter data create(v[:N])
          }

          class A
          {
          public:
          int N;
          B b;
          double *__restrict__ w;
          void construct(int);
          };

          void A::construct(int _N)
          {
          N=_N;
          w=new double[N];
          #pragma acc enter data create(this)
          #pragma acc update device(this)
          #pragma acc enter data create(w[:N])

          // allocate class A
          b.construct(N);
          }


          int main() {

          A myA;
          int N=32;
          myA.construct(N);

          #pragma acc parallel loop present(myA)
          for (int i=0; i<N; ++i) {
          myA.w[i] = i;
          myA.b.u[i] = i;
          myA.b.v[i] = i;
          }
          #pragma acc update host( myA.w[:N], myA.b.u[:N], myA.b.v[:N])
          for (int i=0; i<N; ++i) {
          std::cout << myA.w[i] << ":" << myA.b.u[i] << ":" << myA.b.v[i] << std::endl;
          }
          return 0;
          }
          % pgc++ test.cpp -Minfo=accel -V18.10 -ta=tesla; a.out
          main:
          49, Generating present(myA)
          Accelerator kernel generated
          Generating Tesla code
          52, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
          56, Generating update self(myA.b.u[:N],myA.w[:N],myA.b.v[:N])
          B::construct(int):
          21, Generating update device(this[:1])
          Generating enter data create(this[:1],v[:N],u[:N])
          A::construct(int):
          41, Generating update device(this[:1])
          Generating enter data create(w[:N],this[:1])
          0:0:0
          1:1:1
          2:2:2
          3:3:3
          4:4:4
          5:5:5
          6:6:6
          7:7:7
          8:8:8
          9:9:9
          10:10:10
          11:11:11
          12:12:12
          13:13:13
          14:14:14
          15:15:15
          16:16:16
          17:17:17
          18:18:18
          19:19:19
          20:20:20
          21:21:21
          22:22:22
          23:23:23
          24:24:24
          25:25:25
          26:26:26
          27:27:27
          28:28:28
          29:29:29
          30:30:30
          31:31:31





          share|improve this answer





















          • yes, Can I send the code to you ? Thanks
            – BatiCode
            Dec 27 '18 at 23:41










          • Sure. Send it to PGI Customer Service (trs@pgroup,com) and ask Alex to forward it to me.
            – Mat Colgrove
            Dec 28 '18 at 15:25










          • Thank you so much as I have been struggling with why this situation arises, as both classes are developed separately and they work fine. when I put them together the issue arises. I just submitted it. Thanks again for your time and help
            – BatiCode
            Dec 28 '18 at 16:55














          0












          0








          0






          Do you have a full example which recreates the error you're seeing? I wrote a little test example using your code snip-it and it worked fine. (See below)



          If you were updating the "this" pointer after creating the arrays, then it would be a problem since you'd be overwriting the device pointers with the host pointers. But as you show above, it shouldn't be an issue.



          % cat test.cpp
          #include <iostream>

          class B
          {
          public:
          int N;
          double *__restrict__ u;
          double *__restrict__ v;
          void construct(int);
          };

          void B::construct(int _N)
          {
          N=_N;
          u=new double[N];
          v=new double[N];
          #pragma acc enter data create(this)
          #pragma acc update device(this)
          #pragma acc enter data create(u[:N])
          #pragma acc enter data create(v[:N])
          }

          class A
          {
          public:
          int N;
          B b;
          double *__restrict__ w;
          void construct(int);
          };

          void A::construct(int _N)
          {
          N=_N;
          w=new double[N];
          #pragma acc enter data create(this)
          #pragma acc update device(this)
          #pragma acc enter data create(w[:N])

          // allocate class A
          b.construct(N);
          }


          int main() {

          A myA;
          int N=32;
          myA.construct(N);

          #pragma acc parallel loop present(myA)
          for (int i=0; i<N; ++i) {
          myA.w[i] = i;
          myA.b.u[i] = i;
          myA.b.v[i] = i;
          }
          #pragma acc update host( myA.w[:N], myA.b.u[:N], myA.b.v[:N])
          for (int i=0; i<N; ++i) {
          std::cout << myA.w[i] << ":" << myA.b.u[i] << ":" << myA.b.v[i] << std::endl;
          }
          return 0;
          }
          % pgc++ test.cpp -Minfo=accel -V18.10 -ta=tesla; a.out
          main:
          49, Generating present(myA)
          Accelerator kernel generated
          Generating Tesla code
          52, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
          56, Generating update self(myA.b.u[:N],myA.w[:N],myA.b.v[:N])
          B::construct(int):
          21, Generating update device(this[:1])
          Generating enter data create(this[:1],v[:N],u[:N])
          A::construct(int):
          41, Generating update device(this[:1])
          Generating enter data create(w[:N],this[:1])
          0:0:0
          1:1:1
          2:2:2
          3:3:3
          4:4:4
          5:5:5
          6:6:6
          7:7:7
          8:8:8
          9:9:9
          10:10:10
          11:11:11
          12:12:12
          13:13:13
          14:14:14
          15:15:15
          16:16:16
          17:17:17
          18:18:18
          19:19:19
          20:20:20
          21:21:21
          22:22:22
          23:23:23
          24:24:24
          25:25:25
          26:26:26
          27:27:27
          28:28:28
          29:29:29
          30:30:30
          31:31:31





          share|improve this answer












          Do you have a full example which recreates the error you're seeing? I wrote a little test example using your code snip-it and it worked fine. (See below)



          If you were updating the "this" pointer after creating the arrays, then it would be a problem since you'd be overwriting the device pointers with the host pointers. But as you show above, it shouldn't be an issue.



          % cat test.cpp
          #include <iostream>

          class B
          {
          public:
          int N;
          double *__restrict__ u;
          double *__restrict__ v;
          void construct(int);
          };

          void B::construct(int _N)
          {
          N=_N;
          u=new double[N];
          v=new double[N];
          #pragma acc enter data create(this)
          #pragma acc update device(this)
          #pragma acc enter data create(u[:N])
          #pragma acc enter data create(v[:N])
          }

          class A
          {
          public:
          int N;
          B b;
          double *__restrict__ w;
          void construct(int);
          };

          void A::construct(int _N)
          {
          N=_N;
          w=new double[N];
          #pragma acc enter data create(this)
          #pragma acc update device(this)
          #pragma acc enter data create(w[:N])

          // allocate class A
          b.construct(N);
          }


          int main() {

          A myA;
          int N=32;
          myA.construct(N);

          #pragma acc parallel loop present(myA)
          for (int i=0; i<N; ++i) {
          myA.w[i] = i;
          myA.b.u[i] = i;
          myA.b.v[i] = i;
          }
          #pragma acc update host( myA.w[:N], myA.b.u[:N], myA.b.v[:N])
          for (int i=0; i<N; ++i) {
          std::cout << myA.w[i] << ":" << myA.b.u[i] << ":" << myA.b.v[i] << std::endl;
          }
          return 0;
          }
          % pgc++ test.cpp -Minfo=accel -V18.10 -ta=tesla; a.out
          main:
          49, Generating present(myA)
          Accelerator kernel generated
          Generating Tesla code
          52, #pragma acc loop gang, vector(32) /* blockIdx.x threadIdx.x */
          56, Generating update self(myA.b.u[:N],myA.w[:N],myA.b.v[:N])
          B::construct(int):
          21, Generating update device(this[:1])
          Generating enter data create(this[:1],v[:N],u[:N])
          A::construct(int):
          41, Generating update device(this[:1])
          Generating enter data create(w[:N],this[:1])
          0:0:0
          1:1:1
          2:2:2
          3:3:3
          4:4:4
          5:5:5
          6:6:6
          7:7:7
          8:8:8
          9:9:9
          10:10:10
          11:11:11
          12:12:12
          13:13:13
          14:14:14
          15:15:15
          16:16:16
          17:17:17
          18:18:18
          19:19:19
          20:20:20
          21:21:21
          22:22:22
          23:23:23
          24:24:24
          25:25:25
          26:26:26
          27:27:27
          28:28:28
          29:29:29
          30:30:30
          31:31:31






          share|improve this answer












          share|improve this answer



          share|improve this answer










          answered Dec 27 '18 at 22:54









          Mat Colgrove

          2,17649




          2,17649












          • yes, Can I send the code to you ? Thanks
            – BatiCode
            Dec 27 '18 at 23:41










          • Sure. Send it to PGI Customer Service (trs@pgroup,com) and ask Alex to forward it to me.
            – Mat Colgrove
            Dec 28 '18 at 15:25










          • Thank you so much as I have been struggling with why this situation arises, as both classes are developed separately and they work fine. when I put them together the issue arises. I just submitted it. Thanks again for your time and help
            – BatiCode
            Dec 28 '18 at 16:55


















          • yes, Can I send the code to you ? Thanks
            – BatiCode
            Dec 27 '18 at 23:41










          • Sure. Send it to PGI Customer Service (trs@pgroup,com) and ask Alex to forward it to me.
            – Mat Colgrove
            Dec 28 '18 at 15:25










          • Thank you so much as I have been struggling with why this situation arises, as both classes are developed separately and they work fine. when I put them together the issue arises. I just submitted it. Thanks again for your time and help
            – BatiCode
            Dec 28 '18 at 16:55
















          yes, Can I send the code to you ? Thanks
          – BatiCode
          Dec 27 '18 at 23:41




          yes, Can I send the code to you ? Thanks
          – BatiCode
          Dec 27 '18 at 23:41












          Sure. Send it to PGI Customer Service (trs@pgroup,com) and ask Alex to forward it to me.
          – Mat Colgrove
          Dec 28 '18 at 15:25




          Sure. Send it to PGI Customer Service (trs@pgroup,com) and ask Alex to forward it to me.
          – Mat Colgrove
          Dec 28 '18 at 15:25












          Thank you so much as I have been struggling with why this situation arises, as both classes are developed separately and they work fine. when I put them together the issue arises. I just submitted it. Thanks again for your time and help
          – BatiCode
          Dec 28 '18 at 16:55




          Thank you so much as I have been struggling with why this situation arises, as both classes are developed separately and they work fine. when I put them together the issue arises. I just submitted it. Thanks again for your time and help
          – BatiCode
          Dec 28 '18 at 16:55


















          draft saved

          draft discarded




















































          Thanks for contributing an answer to Stack Overflow!


          • Please be sure to answer the question. Provide details and share your research!

          But avoid



          • Asking for help, clarification, or responding to other answers.

          • Making statements based on opinion; back them up with references or personal experience.


          To learn more, see our tips on writing great answers.





          Some of your past answers have not been well-received, and you're in danger of being blocked from answering.


          Please pay close attention to the following guidance:


          • Please be sure to answer the question. Provide details and share your research!

          But avoid



          • Asking for help, clarification, or responding to other answers.

          • Making statements based on opinion; back them up with references or personal experience.


          To learn more, see our tips on writing great answers.




          draft saved


          draft discarded














          StackExchange.ready(
          function () {
          StackExchange.openid.initPostLogin('.new-post-login', 'https%3a%2f%2fstackoverflow.com%2fquestions%2f53947984%2fnested-classes-with-pointers-woth-openacc%23new-answer', 'question_page');
          }
          );

          Post as a guest















          Required, but never shown





















































          Required, but never shown














          Required, but never shown












          Required, but never shown







          Required, but never shown

































          Required, but never shown














          Required, but never shown












          Required, but never shown







          Required, but never shown







          Popular posts from this blog

          Monofisismo

          Angular Downloading a file using contenturl with Basic Authentication

          Olmecas